Skip to content

Commit 00cb97c

Browse files
author
Ryan M. Bergmann
committed
initial commit of remapping code
1 parent 805d8ba commit 00cb97c

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

62 files changed

+8945
-0
lines changed

‎LCRNG.cuh

+21
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,21 @@
1+
inline __device__ float get_rand(unsigned* in)
2+
{
3+
/*
4+
increments the random number with LCRNG
5+
adapated from OpenMC again
6+
values from http://www.ams.org/journals/mcom/1999-68-225/S0025-5718-99-00996-5/S0025-5718-99-00996-5.pdf
7+
since 32-bit math is being used, 30 bits are used here
8+
*/
9+
const unsigned a = 116646453; // multiplier
10+
const unsigned c = 7; // constant add, must be odd
11+
const unsigned mask = 1073741823; // 2^30-1
12+
const float norm = 9.31322574615478515625E-10; // 2^-30
13+
unsigned nextint = (a * in[0] + c) & mask; // mod by truncation
14+
float randout = nextint*norm;
15+
if(randout>=1.0){
16+
randout=0.9999999;
17+
//printf("RN=1.0 %u %u %10.8E\n",in[0],nextint,randout);
18+
}
19+
in[0]=nextint;
20+
return randout; // return normalized float
21+
}

‎Makefile

+203
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,203 @@
1+
#
2+
CC = gcc
3+
CXX = g++
4+
OPTIX = /usr/local/OptiX-3.0.1
5+
NVCC = nvcc
6+
ARCH = -arch sm_30
7+
C_FLAGS = -O3 -m64 -fPIC
8+
NVCC_FLAGS = -m64 -use_fast_math --compiler-options '-fPIC'
9+
CURAND_LIBS = -lcurand
10+
OPTIX_FLAGS = -I$(OPTIX)/include -L$(OPTIX)/lib64
11+
OPTIX_LIBS = -loptix
12+
CUDA_FLAGS = -I/usr/local/cuda/include -L/usr/local/cuda/lib
13+
CUDPP_PATH = /usr/local/cudpp-2.1/
14+
CUDPP_FLAGS = -I/$(CUDPP_PATH)/include -L/$(CUDPP_PATH)/lib
15+
CUDPP_LIBS = -lcudpp_hash -lcudpp
16+
PYTHON_FLAGS = -I/usr/local/anaconda/include/python2.7 -L/usr/local/anaconda/lib/
17+
PYTHON_LIBS = -lpython2.7
18+
PNG_FLAGS = -L/
19+
PNG_LIBS = -lpng15
20+
LIBS =
21+
22+
23+
COBJS = mt19937ar.o \
24+
print_banner.o \
25+
set_positions_rand.o \
26+
copy_points.o \
27+
macroscopic.o \
28+
microscopic.o \
29+
find_E_grid_index.o \
30+
find_E_grid_index_quad.o \
31+
sample_fission_spectra.o \
32+
sample_fixed_source.o \
33+
tally_spec.o \
34+
escatter.o \
35+
iscatter.o \
36+
cscatter.o \
37+
fission.o \
38+
absorb.o \
39+
make_mask.o \
40+
print_histories.o \
41+
pop_secondaries.o \
42+
pop_source.o \
43+
rebase_yield.o \
44+
flip_done.o \
45+
whistory.o \
46+
wgeometry.o \
47+
optix_stuff.o \
48+
primitive.o \
49+
write_to_file.o \
50+
reaction_edges.o \
51+
device_copies.o \
52+
53+
ptx_objects = camera.ptx \
54+
hits.ptx \
55+
miss.ptx \
56+
box.ptx \
57+
cylinder.ptx \
58+
hex.ptx\
59+
hits_mesh.ptx \
60+
box_mesh.ptx \
61+
cylinder_mesh.ptx \
62+
hex_mesh.ptx\
63+
sphere_mesh.ptx\
64+
65+
66+
all: $(ptx_objects) \
67+
$(COBJS) \
68+
libwarp.so
69+
70+
clean:
71+
rm -f *.ptx *.o *.so gpu debug optixtest
72+
73+
camera.ptx:
74+
$(NVCC) $(ARCH) $(NVCC_FLAGS) $(OPTIX_FLAGS) $(OPTIX_LIBS) -ptx camera.cu
75+
76+
hits.ptx:
77+
$(NVCC) $(ARCH) $(NVCC_FLAGS) $(OPTIX_FLAGS) $(OPTIX_LIBS) -ptx hits.cu
78+
79+
miss.ptx:
80+
$(NVCC) $(ARCH) $(NVCC_FLAGS) $(OPTIX_FLAGS) $(OPTIX_LIBS) -ptx miss.cu
81+
82+
box.ptx:
83+
$(NVCC) $(ARCH) $(NVCC_FLAGS) $(OPTIX_FLAGS) $(OPTIX_LIBS) -ptx box.cu
84+
85+
cylinder.ptx:
86+
$(NVCC) $(ARCH) $(NVCC_FLAGS) $(OPTIX_FLAGS) $(OPTIX_LIBS) -ptx cylinder.cu
87+
88+
hex.ptx:
89+
$(NVCC) $(ARCH) $(NVCC_FLAGS) $(OPTIX_FLAGS) $(OPTIX_LIBS) -ptx hex.cu
90+
91+
hits_mesh.ptx:
92+
$(NVCC) $(ARCH) $(NVCC_FLAGS) $(OPTIX_FLAGS) $(OPTIX_LIBS) -ptx hits_mesh.cu
93+
94+
box_mesh.ptx:
95+
$(NVCC) $(ARCH) $(NVCC_FLAGS) $(OPTIX_FLAGS) $(OPTIX_LIBS) -ptx box_mesh.cu
96+
97+
cylinder_mesh.ptx:
98+
$(NVCC) $(ARCH) $(NVCC_FLAGS) $(OPTIX_FLAGS) $(OPTIX_LIBS) -ptx cylinder_mesh.cu
99+
100+
hex_mesh.ptx:
101+
$(NVCC) $(ARCH) $(NVCC_FLAGS) $(OPTIX_FLAGS) $(OPTIX_LIBS) -ptx hex_mesh.cu
102+
103+
sphere_mesh.ptx:
104+
$(NVCC) $(ARCH) $(NVCC_FLAGS) $(OPTIX_FLAGS) $(OPTIX_LIBS) -ptx sphere_mesh.cu
105+
106+
mt19937ar.o:
107+
$(CXX) $(C_FLAGS) -c -O mt19937ar.cpp
108+
109+
print_banner.o:
110+
$(NVCC) $(NVCC_FLAGS) -c -O print_banner.cpp
111+
112+
set_positions_rand.o:
113+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c set_positions_rand.cu
114+
115+
find_E_grid_index.o:
116+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c find_E_grid_index.cu
117+
118+
find_E_grid_index_quad.o:
119+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c find_E_grid_index_quad.cu
120+
121+
sample_fission_spectra.o:
122+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c sample_fission_spectra.cu
123+
124+
sample_fixed_source.o:
125+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c sample_fixed_source.cu
126+
127+
macroscopic.o:
128+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c macroscopic.cu
129+
130+
microscopic.o:
131+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c microscopic.cu
132+
133+
copy_points.o:
134+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c copy_points.cu
135+
136+
tally_spec.o:
137+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c tally_spec.cu
138+
139+
escatter.o:
140+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c escatter.cu
141+
142+
iscatter.o:
143+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c iscatter.cu
144+
145+
cscatter.o:
146+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c cscatter.cu
147+
148+
fission.o:
149+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c fission.cu
150+
151+
absorb.o:
152+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c absorb.cu
153+
154+
make_mask.o:
155+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c make_mask.cu
156+
157+
print_histories.o:
158+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c print_histories.cu
159+
160+
pop_secondaries.o:
161+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c pop_secondaries.cu
162+
163+
pop_source.o:
164+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c pop_source.cu
165+
166+
rebase_yield.o:
167+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c rebase_yield.cu
168+
169+
flip_done.o:
170+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c flip_done.cu
171+
172+
device_copies.o:
173+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c device_copies.cu
174+
175+
reaction_edges.o:
176+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c reaction_edges.cu
177+
178+
write_to_file.o:
179+
$(NVCC) $(ARCH) $(NVCC_FLAGS) -c write_to_file.cu
180+
181+
whistory.o:
182+
$(CXX) $(C_FLAGS) $(OPTIX_FLAGS) $(CUDPP_FLAGS) $(PNG_FLAGS) $(PYTHON_FLAGS) $(CUDA_FLAGS) -c whistory.cpp
183+
184+
wgeometry.o:
185+
$(CXX) $(C_FLAGS) -c wgeometry.cpp
186+
187+
primitive.o:
188+
$(CXX) $(C_FLAGS) -c primitive.cpp
189+
190+
optix_stuff.o: device_copies.o
191+
$(CXX) -m64 -fPIC $(OPTIX_FLAGS) $(CUDA_FLAGS) $(PNG_FLAGS) -c optix_stuff.cpp
192+
193+
libwarp.so: $(ptx_objects) $(COBJS)
194+
$(NVCC) --shared $(NVCC_FLAGS) $(OPTIX_FLAGS) $(CUDPP_FLAGS) $(PYTHON_FLAGS) $(PNG_FLAGS) $(CURAND_LIBS) $(OPTIX_LIBS) $(CUDPP_LIBS) $(PYTHON_LIBS) $(PNG_LIBS) $(COBJS) -o libwarp.so
195+
196+
gpu: libwarp.so
197+
$(NVCC) -m64 $(OPTIX_FLAGS) $(CUDPP_FLAGS) $(PYTHON_FLAGS) -L/mnt/scratch/gpu-cpp -lwarp main.cpp -o $@
198+
199+
optixtest: libwarp.so
200+
$(NVCC) -m64 $(OPTIX_FLAGS) $(CUDPP_FLAGS) $(PYTHON_FLAGS) -L/Users/rmb/code/gpu-cpp -lwarp -loptix optixtest.cpp -o $@
201+
202+
warp.py: libwarp.so
203+
swig warp.i

‎absorb.cu

+34
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,34 @@
1+
#include <cuda.h>
2+
#include <stdio.h>
3+
#include "datadef.h"
4+
5+
__global__ void absorb_kernel(unsigned N, unsigned* active, unsigned * rxn , unsigned* done){
6+
7+
8+
//PLACEHOLDER FOR FISSIONS, NEED TO READ NU TABLES LATER
9+
10+
int tid = threadIdx.x+blockIdx.x*blockDim.x;
11+
if (tid >= N){return;} //return if out of bounds
12+
13+
//remap to active
14+
//tid=active[tid];
15+
if(done[tid]){return;}
16+
17+
if (rxn[tid] < 102 | rxn[tid] > 800 ){return;} //return if not some sort of absorption, ie (n,not-n)
18+
19+
//printf("in abs, rxn=%u\n",rxn[tid]);
20+
21+
done[tid] = 1;
22+
23+
}
24+
25+
void absorb( cudaStream_t stream, unsigned NUM_THREADS, unsigned N, unsigned* active, unsigned * rxn , unsigned* done){
26+
27+
unsigned blks = ( N + NUM_THREADS - 1 ) / NUM_THREADS;
28+
29+
//absorb_kernel <<< blks, NUM_THREADS >>> ( N, active, rxn , done);
30+
absorb_kernel <<< blks, NUM_THREADS , 0 , stream >>> ( N, active, rxn , done);
31+
cudaThreadSynchronize();
32+
33+
}
34+

‎binary_search.h

+35
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,35 @@
1+
__forceinline__ __device__ unsigned binary_search( float * array , float value, unsigned len ){
2+
3+
// load data
4+
unsigned donesearching = 0;
5+
unsigned cnt = 1;
6+
unsigned powtwo = 2;
7+
int dex = (len) / 2; //N_energiesgth starts at 1, duh
8+
9+
while(!donesearching){
10+
11+
powtwo = powtwo * 2;
12+
if ( array[dex] <= value &&
13+
array[dex+1] > value ) { donesearching = 1; }
14+
else if ( array[dex] > value ) { dex = dex - (( len / powtwo) + 1) ; cnt++; } // +1's are to do a ceiling instead of a floor on integer division
15+
else if ( array[dex] < value ) { dex = dex + (( len / powtwo) + 1) ; cnt++; }
16+
17+
if(cnt>30){
18+
donesearching=1;
19+
printf("device binary search iteration overflow! %p %d val % 10.8f ends % 10.8f % 10.8f\n",array,len,value,array[0],array[len-1]);
20+
dex=0;
21+
}
22+
23+
// edge checks... fix later???
24+
if(dex<0){
25+
dex=0;
26+
}
27+
if(dex>=len){
28+
dex=len-1;
29+
}
30+
}
31+
32+
// output index
33+
return dex;
34+
35+
}

‎box.cu

+39
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,39 @@
1+
#include <optix.h>
2+
#include <optixu/optixu_math_namespace.h>
3+
#include <optixu/optixu_matrix_namespace.h>
4+
#include <optixu/optixu_aabb_namespace.h>
5+
6+
using namespace optix;
7+
8+
rtDeclareVariable(float3, mins, , );
9+
rtDeclareVariable(float3, maxs, , );
10+
rtDeclareVariable(optix::Ray, ray, rtCurrentRay, );
11+
12+
RT_PROGRAM void intersect(int)
13+
{
14+
float3 t0 = (mins - ray.origin)/ray.direction;
15+
float3 t1 = (maxs - ray.origin)/ray.direction;
16+
float3 near = fminf(t0, t1);
17+
float3 far = fmaxf(t0, t1);
18+
float tmin = fmaxf( near );
19+
float tmax = fminf( far );
20+
21+
if(tmin <= tmax) {
22+
bool check_second = true;
23+
if( rtPotentialIntersection( tmin ) ) {
24+
if(rtReportIntersection(0))
25+
check_second = false;
26+
}
27+
if(check_second) {
28+
if( rtPotentialIntersection( tmax ) ) {
29+
rtReportIntersection(0);
30+
}
31+
}
32+
}
33+
}
34+
35+
RT_PROGRAM void bounds (int, float result[6])
36+
{
37+
optix::Aabb* aabb = (optix::Aabb*)result;
38+
aabb->set(mins, maxs);
39+
}

0 commit comments

Comments
 (0)