diff options
Diffstat (limited to 'chroma/cuda/tools.cu')
| -rw-r--r-- | chroma/cuda/tools.cu | 21 |
1 files changed, 21 insertions, 0 deletions
diff --git a/chroma/cuda/tools.cu b/chroma/cuda/tools.cu new file mode 100644 index 0000000..80d06ec --- /dev/null +++ b/chroma/cuda/tools.cu @@ -0,0 +1,21 @@ +//--*-c-*- + +extern "C" +{ + +__global__ void +interleave(int nthreads, unsigned long long *x, unsigned long long *y, + unsigned long long *z, int bits, unsigned long long *dest) +{ + int id = blockIdx.x*blockDim.x + threadIdx.x; + + if (id >= nthreads) + return; + + for (int i=0; i < bits; i++) + dest[id] |= (x[id] & 1 << i) << (2*i) | + (y[id] & 1 << i) << (2*i+1) | + (z[id] & 1 << i) << (2*i+2); +} + +} // extern "C" |