0_Simple__simpleAssert + 0_Simple__simpleAssert_nvrtc

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

▶ 源代码：静态使用

#include <windows.h>
#include <stdio.h>
#include <cuda_runtime.h>
#include "device_launch_parameters.h"
#include <helper_functions.h>
#include <helper_cuda.h>

#define WINDOWS_LEAN_AND_MEAN
#define NOMINMAX

__global__ void testKernel(int N)
{
    int tid = blockIdx.x*blockDim.x + threadIdx.x ;
    
    // 检查条件为“线程总编号小于 N”，即阻塞不满足该条件的线程
    // 阻塞的同时向屏幕输出阻塞线程的信息，包括核函数所在文件绝对路径、行号、线程块号，线程号，没有通过的检查条件
    assert(tid < N) ;
}

bool runTest()
{
    // 使用2个线程块各32条线程，使用 assert() 阻塞最后 4 条（即第 1 线程块的第 29、30、31、32 号线程）
    int Nblocks = 2;
    int Nthreads = 32;
    cudaError_t error ;

    dim3 dimGrid(Nblocks);
    dim3 dimBlock(Nthreads);
    testKernel<<<dimGrid, dimBlock>>>(60);

    printf("
-- Begin assert output

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

");
    
    if (error == cudaErrorAssert)           // 输出错误信息种类
        printf("CUDA error message is: %s
",cudaGetErrorString(error));

    return error == cudaErrorAssert;
}

int main()
{
    bool testResult; 
    
    printf("
	Started!
");

    testResult = runTest();

    printf("
	Completed! main function returned %s
", testResult ? "OK!" : "ERROR!");
    getchar();

    return 0;
}

即时编译版：

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

/*simpleAssert.cpp*/
#include <windows.h>
#include <stdio.h>
#include <cuda_runtime.h>
#include <helper_functions.h>
#include "nvrtc_helper.h"

#define WINDOWS_LEAN_AND_MEAN
#define NOMINMAX

bool runTest()
{
    int Nblocks = 2;
    int Nthreads = 32;

    // 紧张的 .cu 即时编译过程
    char *kernel_file = sdkFindFilePath("simpleAssert_kernel.cu", NULL);

    char *ptx;
    size_t ptxSize;
    compileFileToPTX(kernel_file, 0, NULL, &ptx, &ptxSize);

    CUmodule module = loadPTX(ptx, 1, NULL);
    
    CUfunction kernel_addr;
    cuModuleGetFunction(&kernel_addr, module, "testKernel");

    dim3 dimGrid(Nblocks);
    dim3 dimBlock(Nthreads);
    int count = 60;
    void *args[] = { (void *)&count };
    cuLaunchKernel(kernel_addr,dimGrid.x, dimGrid.y, dimGrid.z,dimBlock.x, dimBlock.y, dimBlock.z,0,0,&args[0],0);

    printf("
-- Begin assert output

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

");

    if (res == CUDA_ERROR_ASSERT)
        printf("Device assert failed as expected
");

    return res == CUDA_ERROR_ASSERT ;
}

int main()
{
    bool testResult;

    printf("
	Started!
");

    testResult = runTest();

    printf("
	Completed! main function returned %s
", testResult ? "OK!" : "ERROR!");
    getchar();

    return 0;
}

▶ 输出结果：

    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。使用的参数与 <<< >>> 方式基本相同。

CUresult CUDAAPI cuLaunchKernel
(
    CUfunction f,
    unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ,
    unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ,
    unsigned int sharedMemBytes,
    CUstream hStream,
    void **kernelParams,
    void **extra
);

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

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

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

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