在核函数中使用强制终止函数 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 = ;
int Nthreads = ;
cudaError_t error ; dim3 dimGrid(Nblocks);
dim3 dimBlock(Nthreads);
testKernel<<<dimGrid, dimBlock>>>(); printf("\n-- Begin assert output\n\n");
error = cudaDeviceSynchronize(); // 使用设备同步来获取错误信息
printf("\n-- End assert output\n\n"); if (error == cudaErrorAssert) // 输出错误信息种类
printf("CUDA error message is: %s\n",cudaGetErrorString(error)); return error == cudaErrorAssert;
} int main()
{
bool testResult; printf("\n\tStarted!\n"); testResult = runTest(); printf("\n\tCompleted! main function returned %s\n", testResult ? "OK!" : "ERROR!");
getchar(); return ;
}
即时编译版：
/*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 = ;
int Nthreads = ; // 紧张的 .cu 即时编译过程
char *kernel_file = sdkFindFilePath("simpleAssert_kernel.cu", NULL); char *ptx;
size_t ptxSize;
compileFileToPTX(kernel_file, , NULL, &ptx, &ptxSize); CUmodule module = loadPTX(ptx, , NULL); CUfunction kernel_addr;
cuModuleGetFunction(&kernel_addr, module, "testKernel"); dim3 dimGrid(Nblocks);
dim3 dimBlock(Nthreads);
int count = ;
void *args[] = { (void *)&count };
cuLaunchKernel(kernel_addr,dimGrid.x, dimGrid.y, dimGrid.z,dimBlock.x, dimBlock.y, dimBlock.z,,,&args[],); printf("\n-- Begin assert output\n\n");
CUresult res = cuCtxSynchronize(); // 用的是上下文同步？
printf("\n-- End assert output\n\n"); if (res == CUDA_ERROR_ASSERT)
printf("Device assert failed as expected\n"); return res == CUDA_ERROR_ASSERT ;
} int main()
{
bool testResult; printf("\n\tStarted!\n"); testResult = runTest(); printf("\n\tCompleted! main function returned %s\n", testResult ? "OK!" : "ERROR!");
getchar(); return ;
}
▶ 输出结果：
Started! -- Begin assert output D:/Program/CUDA/Samples/0_Simple/simpleAssert/simpleAssert.cu:: block: [,,], thread: [,,] Assertion `tid < N` failed.
D:/Program/CUDA/Samples/0_Simple/simpleAssert/simpleAssert.cu:: block: [,,], thread: [,,] Assertion `tid < N` failed.
D:/Program/CUDA/Samples/0_Simple/simpleAssert/simpleAssert.cu:: block: [,,], thread: [,,] Assertion `tid < N` failed.
D:/Program/CUDA/Samples/0_Simple/simpleAssert/simpleAssert.cu:: block: [,,], thread: [,,] 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 中
尚不清楚两者的差别，等待填坑。