Add files via upload

Author: YouQixiaowu

src/02-thread-organization/hello1.py

@@ -0,0 +1 @@
+print('Hello World!\n')

src/02-thread-organization/hello2.py

@@ -0,0 +1,13 @@
+import pycuda.autoinit
+import pycuda.driver as drv
+from pycuda.compiler import SourceModule
+
+mod = SourceModule(r"""
+__global__ void hello_from_gpu(void)
+{
+    printf("hello World from the GPU!\n");  
+}
+""")
+
+hello_from_gpu = mod.get_function("hello_from_gpu")
+hello_from_gpu(grid=(1,1,1),block=(1,1,1))

src/02-thread-organization/hello3.py

@@ -0,0 +1,13 @@
+import pycuda.autoinit
+import pycuda.driver as drv
+from pycuda.compiler import SourceModule
+
+mod = SourceModule(r"""
+__global__ void hello_from_gpu(void)
+{
+    printf("hello World from the GPU!\n");  
+}
+""")
+
+hello_from_gpu = mod.get_function("hello_from_gpu")
+hello_from_gpu(grid=(2,1,1),block=(4,1,1))

src/02-thread-organization/hello4.py

@@ -0,0 +1,15 @@
+import pycuda.autoinit
+import pycuda.driver as drv
+from pycuda.compiler import SourceModule
+
+mod = SourceModule(r"""
+__global__ void hello_from_gpu(void)
+{
+    const int bid = blockIdx.x;
+    const int tid = threadIdx.x;
+    printf("Hello World from block %d and thread %d!\n", bid, tid);
+}
+""")
+
+hello_from_gpu = mod.get_function("hello_from_gpu")
+hello_from_gpu(grid=(2,1,1), block=(4,1,1))

src/02-thread-organization/hello5.py

@@ -0,0 +1,16 @@
+import pycuda.autoinit
+import pycuda.driver as drv
+from pycuda.compiler import SourceModule
+
+mod = SourceModule(r"""
+__global__ void hello_from_gpu(void)
+{
+    const int b = blockIdx.x;
+    const int tx = threadIdx.x;
+    const int ty = threadIdx.y;
+    printf("Hello World from block-%d and thread-(%d, %d)!\n", b, tx, ty);
+}
+""")
+
+hello_from_gpu = mod.get_function("hello_from_gpu")
+hello_from_gpu(grid=(1,1,1), block=(2,4,1))

src/03-basic-framework/add.py

@@ -0,0 +1,11 @@
+import numpy
+
+EPSILON = 1e-15
+a = 1.23
+b = 2.34
+c = 3.57
+N = 100000000
+x = numpy.full((N,1), a)
+y = numpy.full((N,1), b)
+z = x + y
+print('No errors' if (abs(z-c)<EPSILON).all() else 'Has errors')

src/03-basic-framework/add1.py

@@ -0,0 +1,30 @@
+import pycuda.autoinit
+import pycuda.driver as drv
+import numpy
+from pycuda.compiler import SourceModule
+
+mod = SourceModule(r"""
+void __global__ add(const double *x, const double *y, double *z, const int N)
+{
+    const int n = blockDim.x * blockIdx.x + threadIdx.x;
+    z[n] = x[n] + y[n];
+}
+""")
+add = mod.get_function("add")
+
+EPSILON = 1e-15
+a = 1.23
+b = 2.34
+c = 3.57
+N = 100000000
+h_x = numpy.full((N,1), a)
+h_y = numpy.full((N,1), b)
+h_z = numpy.zeros_like(h_x)
+d_x = drv.mem_alloc(h_x.nbytes)
+d_y = drv.mem_alloc(h_y.nbytes)
+d_z = drv.mem_alloc(h_z.nbytes)
+drv.memcpy_htod(d_x, h_x)
+drv.memcpy_htod(d_y, h_y)
+add(d_x, d_y, d_z, grid=(N//128, 1), block=(128,1,1))
+drv.memcpy_dtoh(h_z, d_z)
+print('No errors' if (abs(h_z-c)<EPSILON).all() else 'Has errors')

src/03-basic-framework/add2wrong.py

@@ -0,0 +1,30 @@
+import pycuda.autoinit
+import pycuda.driver as drv
+import numpy
+from pycuda.compiler import SourceModule
+
+mod = SourceModule(r"""
+void __global__ add(const double *x, const double *y, double *z, const int N)
+{
+    const int n = blockDim.x * blockIdx.x + threadIdx.x;
+    z[n] = x[n] + y[n];
+}
+""")
+add = mod.get_function("add")
+
+EPSILON = 1e-15
+a = 1.23
+b = 2.34
+c = 3.57
+N = 100000000
+h_x = numpy.full((N,1), a)
+h_y = numpy.full((N,1), b)
+h_z = numpy.zeros_like(h_x)
+d_x = drv.mem_alloc(h_x.nbytes)
+d_y = drv.mem_alloc(h_y.nbytes)
+d_z = drv.mem_alloc(h_z.nbytes)
+drv.memcpy_dtoh(d_x, h_x)
+drv.memcpy_dtoh(d_y, h_y)
+add(d_x, d_y, d_z, grid=(N//128, 1), block=(128,1,1))
+drv.memcpy_dtoh(h_z, d_z)
+print('No errors' if (abs(h_z-c)<EPSILON).all() else 'Has errors')

src/03-basic-framework/add3if.py

@@ -0,0 +1,33 @@
+import pycuda.autoinit
+import pycuda.driver as drv
+import numpy
+from pycuda.compiler import SourceModule
+
+mod = SourceModule(r"""
+void __global__ add(const double *x, const double *y, double *z, const int N)
+{
+    const int n = blockDim.x * blockIdx.x + threadIdx.x;
+    if(n < N) 
+    {
+        z[n] = x[n] + y[n];
+    }
+}
+""")
+add = mod.get_function("add")
+
+EPSILON = 1e-15
+a = 1.23
+b = 2.34
+c = 3.57
+N = 100000000
+h_x = numpy.full((N,1), a)
+h_y = numpy.full((N,1), b)
+h_z = numpy.zeros_like(h_x)
+d_x = drv.mem_alloc(h_x.nbytes)
+d_y = drv.mem_alloc(h_y.nbytes)
+d_z = drv.mem_alloc(h_z.nbytes)
+drv.memcpy_htod(d_x, h_x)
+drv.memcpy_htod(d_y, h_y)
+add(d_x, d_y, d_z, numpy.int32(N), grid=((N-1)//128+1, 1), block=(128, 1, 1))
+drv.memcpy_dtoh(h_z, d_z)
+print('No errors' if (abs(h_z-c)<EPSILON).all() else 'Has errors')

src/03-basic-framework/add4device.py

@@ -0,0 +1,84 @@
+import pycuda.autoinit
+import pycuda.driver as drv
+import numpy
+from pycuda.compiler import DynamicSourceModule
+
+import numpy
+
+
+mod1 = DynamicSourceModule(r"""
+double __device__ add1_device(double x, double y)
+{
+    return (x + y);
+}
+
+void __global__ add1(double *x, double *y, double *z, int N)
+{
+    const int n = blockDim.x * blockIdx.x + threadIdx.x;
+    if(n < N) 
+    {
+        z[n] = add1_device(x[n], y[n]);
+    }
+}
+""")
+add1 = mod1.get_function("add1")
+
+mod2 = DynamicSourceModule(r"""
+double __device__ add2_device(double x, double y, double *z)
+{
+    *z = x + y;
+}
+
+void __global__ add2(double *x, double *y, double *z, int N)
+{
+    const int n = blockDim.x * blockIdx.x + threadIdx.x;
+    if(n < N) 
+    {
+        add2_device(x[n], y[n], &z[n]);
+    }
+}
+""")
+add2 = mod2.get_function("add2")
+
+mod3 = DynamicSourceModule(r"""
+double __device__ add3_device(double x, double y, double &z)
+{
+    z = x + y;
+}
+
+void __global__ add3(double *x, double *y, double *z, int N)
+{
+    const int n = blockDim.x * blockIdx.x + threadIdx.x;
+    if(n < N) 
+    {
+        add3_device(x[n], y[n], z[n]);
+    }
+}
+""")
+add3 = mod3.get_function("add3")
+
+EPSILON = 1e-15
+a = 1.23
+b = 2.34
+c = 3.57
+N = 100000001
+h_x = numpy.full((N,1), a)
+h_y = numpy.full((N,1), b)
+h_z = numpy.zeros_like(h_x)
+d_x = drv.mem_alloc(h_x.nbytes)
+d_y = drv.mem_alloc(h_y.nbytes)
+d_z = drv.mem_alloc(h_z.nbytes)
+drv.memcpy_htod(d_x, h_x)
+drv.memcpy_htod(d_y, h_y)
+
+add1(d_x, d_y, d_z, numpy.int32(N), grid=((N-1)//128+1, 1), block=(128, 1, 1))
+drv.memcpy_dtoh(h_z, d_z)
+print('No errors' if (abs(h_z-c)<EPSILON).all() else 'Has errors')
+
+add2(d_x, d_y, d_z, numpy.int32(N), grid=((N-1)//128+1, 1), block=(128, 1, 1))
+drv.memcpy_dtoh(h_z, d_z)
+print('No errors' if (abs(h_z-c)<EPSILON).all() else 'Has errors')
+
+add3(d_x, d_y, d_z, numpy.int32(N), grid=((N-1)//128+1, 1), block=(128, 1, 1))
+drv.memcpy_dtoh(h_z, d_z)
+print('No errors' if (abs(h_z-c)<EPSILON).all() else 'Has errors')