|
| 1 | +#include "simphox_gpu.h" |
| 2 | + |
| 3 | +#include <cuda_runtime.h> |
| 4 | +#include <curand_kernel.h> |
| 5 | + |
| 6 | +#include "simphony/sysrap/CUDA_CHECK.h" |
| 7 | +#include "simphony/sysrap/srng.h" |
| 8 | + |
| 9 | +namespace |
| 10 | +{ |
| 11 | + |
| 12 | +__global__ void generate_photons_kernel(storch torch, sphoton* photons, unsigned int num_photons, unsigned int seed) |
| 13 | +{ |
| 14 | + unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x; |
| 15 | + if (idx >= num_photons) |
| 16 | + { |
| 17 | + return; |
| 18 | + } |
| 19 | + |
| 20 | + curandStatePhilox4_32_10 rng; |
| 21 | + curand_init(seed, idx, 0, &rng); |
| 22 | + |
| 23 | + qtorch qt; |
| 24 | + qt.t = torch; |
| 25 | + |
| 26 | + sphoton photon; |
| 27 | + storch::generate(photon, rng, qt.q, idx, 0); |
| 28 | + photons[idx] = photon; |
| 29 | +} |
| 30 | + |
| 31 | +} // namespace |
| 32 | + |
| 33 | +std::vector<sphoton> generate_photons_gpu(const storch& torch, unsigned int num_photons, unsigned int seed) |
| 34 | +{ |
| 35 | + if (num_photons == 0) |
| 36 | + { |
| 37 | + num_photons = torch.numphoton; |
| 38 | + } |
| 39 | + |
| 40 | + std::vector<sphoton> photons(num_photons); |
| 41 | + if (num_photons == 0) |
| 42 | + { |
| 43 | + return photons; |
| 44 | + } |
| 45 | + |
| 46 | + sphoton* d_photons = nullptr; |
| 47 | + CUDA_CHECK(cudaMalloc(reinterpret_cast<void**>(&d_photons), photons.size() * sizeof(sphoton))); |
| 48 | + |
| 49 | + constexpr unsigned int threads_per_block = 256; |
| 50 | + unsigned int blocks = (num_photons + threads_per_block - 1) / threads_per_block; |
| 51 | + |
| 52 | + generate_photons_kernel<<<blocks, threads_per_block>>>(torch, d_photons, num_photons, seed); |
| 53 | + CUDA_SYNC_CHECK(); |
| 54 | + |
| 55 | + CUDA_CHECK(cudaMemcpy(photons.data(), d_photons, photons.size() * sizeof(sphoton), cudaMemcpyDeviceToHost)); |
| 56 | + CUDA_CHECK(cudaFree(d_photons)); |
| 57 | + |
| 58 | + return photons; |
| 59 | +} |
0 commit comments