CUDA编程（二）--基础关键字及函数说明

1.基础关键字

• host（CPU，主机编译、主机运行）

  • __host__定义了主机端的函数，在主机端调用，主机端执行的函数，也就是我们正常的c/c++的函数，如果不加任何的修饰符，默认就是__host__ 函数，这些函数只为主机端编译；

  • __global__和 __host__不能同时使用；

  • 但是__device__和__host__可以同时使用，这个函数会同时为主机端和设备端编译。代码中可以用__CUDA_ARCH__这个宏来区分主机和设备端以及不同架构的代码。

• global（GPU，主机调用GPU执行）

  • __global__ 修饰的函数是核函数，这些函数 在GPU上执行 ，但是需要 在CPU端调用

  • PS.计算能力3.5以上的引入了 CUDA Dynamic Parallelism机制可以在devices端调用核函数；

  • __global__修饰的函数必须采用void返回值，并且需要在调用时制定运行的参数 (也就是<<<>>>里的block数和线程数）；

  • 同时__global__函数是异步的，这也代表着函数没被执行完就返回了控制权，所以测量核函数的时间需要同步操作才能获得准确的结果。

__global__ void MatmulKernel(float *M_device, float *N_device, float *P_device, int width){
    //
    //
}

//核函数调用
MatmulKernel <<<dimGrid, dimBlock>>> (M_device, N_device, P_device, width);

// host和device的同步
cudaDeviceSynchronize();

• device（GPU，GPU调用GPU执行）

  • __device__函数是在GPU端调用且在GPU端执行的的函数，__global___和__device__不能同时使用；

  • __device__在编译器中会被编译为内联函数。

• noinline 和 forceinline

  • __device__会被编译器编译为内联函数，__noinline__和 __forceinline__则指定编译器不要内联，和强制内联这些函数。这两个符号不能同时使用。

2. CUDA core

2.1 CUDA runtime api

• 以cuda*开头的api一般被称作cuda runtime api,以CU*开头的api被称作cuda driver api。 cuda runtime api对CUDA Driver的操作做封装便于使用。包括：
  • implicit initialization (隐式初始化）

  • context management (上下文管理）

  • module management (模块管理）

• 分配内存空间

  • cudaMalloc (在device端分配空间)，是一种cuda runtime api（*）；

  • cudaMallocHost (在host端的pinned memory上分配空间）。

• 数据传输

  • cudaMemcpy (以同步的方式，将数据在host->device, device->device, device->host进行传输）cudaMemcpy(M_device, M_host, size, cudaMemcpyHostToDevice);GPU内存分配成功，函数返回：cudasuccess;失败：cudaErrorMemoryAllocation。可以使用char *cudaGetErrorString(cudaError_t error)将error_t 错误码转换为string可读的错误消息。

  • cudaMemcpyAsync (以异步的方式，进行数据传输）

CUDA Driver api 和 CUDA runtime api

• host与device的数据传输

• CUDA中有个规定，就是一个block中可以分配的thread的数量最大是1,024个线程。如果大于1,024会显示配置错误。

• 第一次使用GPU时，GPU启动和初始化会花费一些时间，此时计算耗时不准确，因此一般在计算耗时工作时，一般先进行warm up，即让GPU先做一些任务然后再去做需要计算的任务。

• cudaDeviceSynchronize():host和device的同步。也就是cpu和gpu执行的同步，因为kernel函数的执行默认是异步的。所以需要cudaDeviceSynchroize()来强制性的让kernel函数的结果执行结束之后host再执行下一步。

• 释放资源

  • cudaDeviceRest（）显式地释放和清空当前进程中与当前设备有关的所有资源。