eic-opticks/qudarap/QCurandStateMonolithic.cu

0001 #include <cstdio>
0002 #include "curand_kernel.h"
0003 #include "qcurandwrap.h"
0004 #include "qrng.h"
0005
0006 #include "QUDA_CHECK.h"
0007 #include "SLaunchSequence.h"
0008
0009
0010 __global__ void _QCurandStateMonolithic_curand_init(int threads_per_launch, int id_offset, qcurandwrap<XORWOW>* cs, XORWOW* states_thread_offset )
0011 {
0012     int id = blockIdx.x*blockDim.x + threadIdx.x;
0013     if (id >= threads_per_launch) return;
0014     curand_init(cs->seed, id+id_offset, cs->offset, states_thread_offset + id );
0015
0016     //if( id == 0 ) printf("// _QCurandStateMonolithic_curand_init id_offset %d \n", id_offset );
0017 }
0018
0019
0020
0021 extern "C" void QCurandStateMonolithic_curand_init(SLaunchSequence* seq,  qcurandwrap<XORWOW>* cs, qcurandwrap<XORWOW>* d_cs)
0022 {
0023     // NB this is still on CPU, dereferencing d_cs here will BUS_ERROR
0024
0025     printf("//QCurandStateMonolithic_curand_init seq.items %d cs %p  d_cs %p cs.num %llu \n", seq->items, cs, d_cs, cs->num );
0026
0027     cudaEvent_t start, stop ;
0028
0029     for(unsigned i=0 ; i < seq->launches.size() ; i++)
0030     {
0031         SLaunch& l = seq->launches[i] ;
0032         printf("// l.sequence_index %d  l.blocks_per_launch %d l.threads_per_block %d  l.threads_per_launch %d l.thread_offset %d  \n",
0033                    l.sequence_index,    l.blocks_per_launch,   l.threads_per_block,    l.threads_per_launch,   l.thread_offset  );
0034
0035         int id_offset = l.thread_offset ;
0036
0037         XORWOW* states_thread_offset = cs->states  + l.thread_offset ;
0038
0039         QUDA::before_kernel( start, stop );
0040
0041         _QCurandStateMonolithic_curand_init<<<l.blocks_per_launch,l.threads_per_block>>>( l.threads_per_launch, id_offset, d_cs, states_thread_offset  );
0042
0043         l.kernel_time = QUDA::after_kernel( start, stop );
0044     }
0045 }
0046
0047