0_Simple__simpleAssert + 0_Simple__simpleAssert

0_Simple__simpleAssert + 0_Simple__simpleAssert_nvrtc

在核函数中使用强制终止函数 assert()。并且在静态代码和运行时编译两种条件下使用。

▶ 源代码：静态使用

 1 #include <windows.h>
 2 #include <stdio.h>
 3 #include <cuda_runtime.h>
 4 #include "device_launch_parameters.h"
 5 #include <helper_functions.h>
 6 #include <helper_cuda.h>
 7 
 8 #define WINDOWS_LEAN_AND_MEAN
 9 #define NOMINMAX
10 
11 __global__ void testKernel(int N)
12 {
13     int tid = blockIdx.x*blockDim.x + threadIdx.x ;
14     
15     // 检查条件为“线程总编号小于 N”，即阻塞不满足该条件的线程
16     // 阻塞的同时向屏幕输出阻塞线程的信息，包括核函数所在文件绝对路径、行号、线程块号，线程号，没有通过的检查条件
17     assert(tid < N) ;
18 }
19 
20 bool runTest()
21 {
22     // 使用2个线程块各32条线程，使用 assert() 阻塞最后 4 条（即第 1 线程块的第 29、30、31、32 号线程）
23     int Nblocks = 2;
24     int Nthreads = 32;
25     cudaError_t error ;
26 
27     dim3 dimGrid(Nblocks);
28     dim3 dimBlock(Nthreads);
29     testKernel<<<dimGrid, dimBlock>>>(60);
30 
31     printf("
-- Begin assert output

");
32     error = cudaDeviceSynchronize();        // 使用设备同步来获取错误信息
33     printf("
-- End assert output

");
34     
35     if (error == cudaErrorAssert)           // 输出错误信息种类
36         printf("CUDA error message is: %s
",cudaGetErrorString(error));
37 
38     return error == cudaErrorAssert;
39 }
40 
41 int main()
42 {
43     bool testResult; 
44     
45     printf("
	Started!
");
46 
47     testResult = runTest();
48 
49     printf("
	Completed! main function returned %s
", testResult ? "OK!" : "ERROR!");
50     getchar();
51 
52     return 0;
53 }

即时编译版：

1 /*simpleAssert_kernel.cu*/
2 extern "C" __global__ void testKernel(int N)
3 {
4     int tid = blockIdx.x*blockDim.x + threadIdx.x ;
5     assert(tid < N) ;
6 }

 1 /*simpleAssert.cpp*/
 2 #include <windows.h>
 3 #include <stdio.h>
 4 #include <cuda_runtime.h>
 5 #include <helper_functions.h>
 6 #include "nvrtc_helper.h"
 7 
 8 #define WINDOWS_LEAN_AND_MEAN
 9 #define NOMINMAX
10 
11 bool runTest()
12 {
13     int Nblocks = 2;
14     int Nthreads = 32;
15 
16     // 紧张的 .cu 即时编译过程
17     char *kernel_file = sdkFindFilePath("simpleAssert_kernel.cu", NULL);
18 
19     char *ptx;
20     size_t ptxSize;
21     compileFileToPTX(kernel_file, 0, NULL, &ptx, &ptxSize);
22 
23     CUmodule module = loadPTX(ptx, 1, NULL);
24     
25     CUfunction kernel_addr;
26     cuModuleGetFunction(&kernel_addr, module, "testKernel");
27 
28     dim3 dimGrid(Nblocks);
29     dim3 dimBlock(Nthreads);
30     int count = 60;
31     void *args[] = { (void *)&count };
32     cuLaunchKernel(kernel_addr,dimGrid.x, dimGrid.y, dimGrid.z,dimBlock.x, dimBlock.y, dimBlock.z,0,0,&args[0],0);
33 
34     printf("
-- Begin assert output

");
35     CUresult res = cuCtxSynchronize();      // 用的是上下文同步？
36     printf("
-- End assert output

");
37 
38     if (res == CUDA_ERROR_ASSERT)
39         printf("Device assert failed as expected
");
40 
41     return res == CUDA_ERROR_ASSERT ;
42 }
43 
44 int main()
45 {
46     bool testResult;
47 
48     printf("
	Started!
");
49 
50     testResult = runTest();
51 
52     printf("
	Completed! main function returned %s
", testResult ? "OK!" : "ERROR!");
53     getchar();
54 
55     return 0;
56 }

▶ 输出结果：

    Started!

-- Begin assert output

D:/Program/CUDA/Samples/0_Simple/simpleAssert/simpleAssert.cu:18: block: [1,0,0], thread: [28,0,0] Assertion `tid < N` failed.
D:/Program/CUDA/Samples/0_Simple/simpleAssert/simpleAssert.cu:18: block: [1,0,0], thread: [29,0,0] Assertion `tid < N` failed.
D:/Program/CUDA/Samples/0_Simple/simpleAssert/simpleAssert.cu:18: block: [1,0,0], thread: [30,0,0] Assertion `tid < N` failed.
D:/Program/CUDA/Samples/0_Simple/simpleAssert/simpleAssert.cu:18: block: [1,0,0], thread: [31,0,0] Assertion `tid < N` failed.

-- End assert output

CUDA error message is: device-side assert triggered

    Completed! main function returned OK!

▶ 涨姿势：

● 在核函数中使用  assert( condition ) 来检查各线程中是否满足某条件。
    若不满足条件 condition，则强制终止该线程，并输出核函数所在文件绝对路径、行号、线程块号，线程号，没有通过的检查条件
    返回错误种类： cudaErrorAssert，错误代码 59，信息为 device-side assert triggered
    cudaErrorAssert 为定义在 driver_type.h 中的枚举类型 enum __device_builtin__ cudaError{...}; 中，记录了各种错误信息。

● 调用核函数的另一种方法。使用定义在 cuda.h 中的函数 cuLaunchKernel。使用的参数与 <<< >>> 方式基本相同。

 1 CUresult CUDAAPI cuLaunchKernel
 2 (
 3     CUfunction f,
 4     unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ,
 5     unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ,
 6     unsigned int sharedMemBytes,
 7     CUstream hStream,
 8     void **kernelParams,
 9     void **extra
10 );

● 两种方法使用的同步函数

　　静态方法时使用的是设备同步 extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaDeviceSynchronize(void);，定义在 cuda_runtime_api.h 中

　　即时编译时用的是上下文同步 CUresult CUDAAPI cuCtxSynchronize(void); ，定义在 cuda.h 中

　　尚不清楚两者的差别，等待填坑。