31__global__
void generate_phase_factors_gpu_kernel
35 double const* atom_pos,
37 acc_complex_double_t* phase_factors
41 int igloc = blockIdx.x * blockDim.x + threadIdx.x;
43 if (igloc < num_gvec_loc) {
44 int gvx = gvec[array2D_offset(igloc, 0, num_gvec_loc)];
45 int gvy = gvec[array2D_offset(igloc, 1, num_gvec_loc)];
46 int gvz = gvec[array2D_offset(igloc, 2, num_gvec_loc)];
48 double ax = atom_pos[array2D_offset(ia, 0, num_atoms)];
49 double ay = atom_pos[array2D_offset(ia, 1, num_atoms)];
50 double az = atom_pos[array2D_offset(ia, 2, num_atoms)];
52 double p = twopi * (ax * gvx + ay * gvy + az * gvz);
57 phase_factors[array2D_offset(igloc, ia, num_gvec_loc)] = make_accDoubleComplex(cosp, sinp);
62extern "C" void generate_phase_factors_gpu(
int num_gvec_loc__,
65 double const* atom_pos__,
66 acc_complex_double_t* phase_factors__)
70 dim3 grid_b(num_blocks(num_gvec_loc__, grid_t.x), num_atoms__);
72 accLaunchKernel((generate_phase_factors_gpu_kernel), dim3(grid_b), dim3(grid_t), 0, 0,
Common device functions used by GPU kernels.
Uniform interface to the runtime API of CUDA and ROCm.
Namespace for accelerator-related functions.
Namespace of the SIRIUS library.