forked from weft/warp
-
Notifications
You must be signed in to change notification settings - Fork 1
/
Copy pathsample_fixed_source.cu
59 lines (43 loc) · 1.63 KB
/
sample_fixed_source.cu
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
#include <cuda.h>
#include <stdio.h>
#include "datadef.h"
#include "LCRNG.cuh"
#include "check_cuda.h"
__global__ void sample_fixed_source_kernel(unsigned N, unsigned* active, unsigned* rn_bank , float * E, source_point* space){
int tid = threadIdx.x+blockIdx.x*blockDim.x;
if (tid >= N){return;}
//remap to active
//tid=active[tid];
// load in
unsigned rn = rn_bank[ tid ];
const float pi = 3.14159265359 ;
const float mu = ( get_rand(&rn) ) * 2.0 - 1.0;
const float theta = ( get_rand(&rn) ) * 2.0 * pi ;
//monoenergetic for now
E[tid]=10.0; //1.0e-6;
//point source for now
space[tid].x = 0.0;
space[tid].y = 0.0;
space[tid].z = 0.0;
//set isotropic for now
space[tid].xhat = sqrtf(1-mu*mu) * cosf( theta );
space[tid].yhat = sqrtf(1-mu*mu) * sinf( theta );
space[tid].zhat = mu;
rn_bank[tid] = rn;
}
/**
* \brief a
* \details b
*
* @param[in] NUM_THREADS - the number of threads to run per thread block
* @param[in] N - the total number of threads to launch on the grid
* @param[in] active - device pointer to remapping vector
* @param[in] rn_bank - device pointer to random number array
* @param[in] E - device pointer to energy data array
* @param[in] space - device pointer to spatial data array
*/
void sample_fixed_source( unsigned NUM_THREADS, unsigned N, unsigned* active, unsigned* rn_bank, float * E, source_point* space){
unsigned blks = ( N + NUM_THREADS - 1 ) / NUM_THREADS;
sample_fixed_source_kernel <<< blks, NUM_THREADS >>> ( N, active, rn_bank, E , space );
check_cuda(cudaThreadSynchronize());
}