summaryrefslogtreecommitdiffstats
path: root/functionTest.cu
diff options
context:
space:
mode:
authorVolker Hoffmann <volker@cheleb.net>2014-07-04 13:49:20 +0200
committerVolker Hoffmann <volker@cheleb.net>2014-07-04 13:49:20 +0200
commitfe66d798f59da25f826c547313b779cee0c9f443 (patch)
tree49f7d6596238760181e630f9dbb91425d2c4ada8 /functionTest.cu
parent36c6e1cf8e9894abaf5d47de0ee5107c0bbec8c0 (diff)
shared memory and device function demo
Diffstat (limited to 'functionTest.cu')
-rw-r--r--functionTest.cu47
1 files changed, 47 insertions, 0 deletions
diff --git a/functionTest.cu b/functionTest.cu
new file mode 100644
index 0000000..d0f066d
--- /dev/null
+++ b/functionTest.cu
@@ -0,0 +1,47 @@
+#include <cuda.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+__device__ void f2(int idx, float x_s[4])
+{
+ if(idx>0) {
+ printf("%i %.2e %.2e\n", idx, x_s[idx], x_s[idx-1]);
+ } else {
+ printf("%i %.2e\n", idx, x_s[idx]);
+ }
+}
+
+__global__ void f1(float *x_d)
+{
+ int idx = blockIdx.x * blockDim.x + threadIdx.x;
+ __shared__ float x_s[4];
+ x_s[idx] = x_d[idx];
+ __syncthreads();
+ //printf("%i %.2e %.2e\n", idx, x_d[idx], x_s[idx]);
+ f2(idx, x_s);
+}
+
+int main()
+{
+ float *x_h, *x_d;
+ int nx = 4;
+
+ // allocate, copy
+ cudaSetDevice(0);
+ x_h = (float *) malloc(sizeof(float)*nx);
+ for(int ii=0; ii<nx; ii++) {
+ x_h[ii] = ii;
+ printf("%i %.2e\n", ii, x_h[ii]);
+ }
+ printf("\n");
+ cudaMalloc((void**) &x_d, sizeof(float)*nx);
+ cudaMemcpy(x_d, x_h, sizeof(float)*nx, cudaMemcpyHostToDevice);
+
+ // run kernel
+ f1 <<< 1, 4 >>> (x_d);
+
+ // free
+ cudaFree(x_d);
+ free(x_h);
+}
+