CUDA samples系列 0.2 simpleAssert

目录

获取相关信息的库

设置核函数的线程数

定义核函数与处理错误信息

结果

扩展：simpleAssert_nvrtc

---

上回发现samples里的排序并不是按照难度等级依次上升的，所以现在就从simple里面依次讲起，因为是simple开头的，想来也是智障等级的例程吧。

获取相关信息的库

#include <sys/utsname.h>

在这个库中，有个结构utsname，里面可以获取这些信息。

struct utsname
  { char sysname[_UTSNAME_SYSNAME_LENGTH];//当前操作系统名
   char nodename[_UTSNAME_NODENAME_LENGTH];//网络上的名称
   char release[_UTSNAME_RELEASE_LENGTH];//当前发布级别
   char version[_UTSNAME_VERSION_LENGTH];//当前发布版本
   char machine[_UTSNAME_MACHINE_LENGTH];//当前硬件体系类型
#if _UTSNAME_DOMAIN_LENGTH - 0
    /* Name of the domain of this node on the network.  */
# ifdef __USE_GNU
    char domainname[_UTSNAME_DOMAIN_LENGTH]; //当前域名
# else
    char __domainname[_UTSNAME_DOMAIN_LENGTH];
# endif
#endif
  };

获取方式类似如下：

    utsname OS_System_Type;
    uname(&OS_System_Type);

    printf("OS_System_Type.release = %s\n", OS_System_Type.release);

然后是获取gpu信息的函数，上回0.1提到过

    // This will pick the best possible CUDA capable device
    devID = findCudaDevice(argc, (const char **)argv);

    checkCudaErrors(cudaGetDeviceProperties(&deviceProp, devID));

    if (deviceProp.major < 2)
    {
        printf("simpleAssert requires a GPU with compute capability "
               "2.0 or later, exiting...\n");

        exit(EXIT_SUCCESS);
    }

设置核函数的线程数

    int Nblocks = 2;
    int Nthreads = 32;    
    dim3 dimGrid(Nblocks);
    dim3 dimBlock(Nthreads);

dim3 这个类型，定义核函数的block个数，thread个数，基本都用这个类型。可以是1,2,3维的，这里是一维的。

2、3维的比如这样：blockSize(5,5,5)， dim3 blockSize(5,4);

实际上，你如果只写一维的话，比如上面代码里dim3 dimGrid(Nblocks)， 它会看作是dim3 dimGrid(Nblocks，1，1)，最后总归是三维的。

定义核函数与处理错误信息

cudaError_t error ;
error = cudaDeviceSynchronize();

先看这句。cudaDeviceSynchronize() 会阻塞当前程序的执行，直到所有任务都处理完毕；也就是说，程序走到这里，会等待这句话之前的所有代码全部执行完毕了，所有的stream都执行完毕了，关于stream，上一篇0.1已经做了详细直观的解释。

也有与stream绑定的阻塞等待函数：cudaStreamSynchronize(streamID)带有一个参数，cuda流ID，它只阻塞那些cuda流ID等于参数中指定ID的那些cuda例程，对于那些流ID不等的例程，还是异步执行的。

这里error用来接收错误信息，也就是这句话之前的程序在执行过程中得到的错误。全部正确，返回0，即“cudaSuccess ”，没有错误，否则返回错误编号。

    if (error == cudaErrorAssert)
    {
        printf("Device assert failed as expected, "
               "CUDA error message is: %s\n\n",
               cudaGetErrorString(error));
    }

这里有个cudaErrorAssert，是个没有定义的，应该是专有名词，最后再cuda官方教程中找到了

在这里：https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html

搜索“cudaErrorAssert”可以看到，应该是说核函数里的断言错误。

__global__ void testKernel(int N)
{
    int gtid = blockIdx.x*blockDim.x + threadIdx.x ;
    assert(gtid < N) ;
}

这是核函数，输入的60，计算出线程号,断言线程号<60后。开辟的线程数2*32=64，线程号从0开始，所以60,61,62,63这4个线程号会被断言错误。

结果

扩展：simpleAssert_nvrtc

这是与simpleAssert实验同样功能的一份代码。

这里的区别是：通过nvtc编译库，来编译cuda核函数，同时，将除此以外的代码全部放在一个cpp文件里去。

至于为什么这么做，可以参考这里：

https://blog.csdn.net/q583956932/article/details/78764301

简单概括下就是：不这么做的话，就得用cuda c/c++编译器来编译cpp代码，这是非常慢的。这样的话，可以用系统自带的c/c++编译器编译cpp，用nvrtc编译cu文件。编译很快。

看了下感觉就这样，源代码注释很详细，没必要详细说了。而且目前暂时一般也不会用到这个nvrtc来分开编译cpp与cuda。