forked from HeLiuCMU/HEXOMAP
-
Notifications
You must be signed in to change notification settings - Fork 0
/
Copy pathRandomCuda.py
51 lines (43 loc) · 1.4 KB
/
RandomCuda.py
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
# Generating Random numbers and random rotation matrix
# Auther
import numpy as np
import pycuda.autoinit
from pycuda.compiler import SourceModule
from pycuda import gpuarray
code = """
#include <curand_kernel.h>
const int nstates = %(NGENERATORS)s;
__device__ curandState_t* states[nstates];
__global__ void initkernel(int seed)
{
int tidx = threadIdx.x + blockIdx.x * blockDim.x;
if (tidx < nstates) {
curandState_t* s = new curandState_t;
if (s != 0) {
curand_init(seed, tidx, 0, s);
}
states[tidx] = s;
}
}
__global__ void randfillkernel(float *values, int N)
{
int tidx = threadIdx.x + blockIdx.x * blockDim.x;
if (tidx < nstates) {
curandState_t s = *states[tidx];
for(int i=tidx; i < N; i += blockDim.x * gridDim.x) {
values[i] = curand_uniform(&s);
}
*states[tidx] = s;
}
}
"""
def random_gpuarray(N,nValues)
#N = 1024
mod = SourceModule(code % { "NGENERATORS" : N }, no_extern_c=True, arch="sm_52")
init_func = mod.get_function("_Z10initkerneli")
fill_func = mod.get_function("_Z14randfillkernelPfi")
seed = np.int32(123456789)
nvalues = 10 * N
init_func(seed, block=(N,1,1), grid=(1,1,1))
gdata = gpuarray.zeros(nvalues, dtype=np.float32)
fill_func(gdata, np.int32(nvalues), block=(N,1,1), grid=(1,1,1))