2017年7月6日 周四

时间：2021-01-12 14:01:34

CUDA学习

CUDA Runtime

初始化

没有显式的初始化runtime的函数，在一个runtime第一次被调用时完成初始化。

初始化时，runtime为系统中的每个device创建一个CUDA context，然后将它们加载到device memory中。当一个host thread调用cudaDeviceReset()时，将销毁该host thread当前操作的device的primary context。之后任何将该device设置为current的host thread会为该device创建一个新的primary context。

Device Memory

在分配device memory时，可以以linear memory的形式或CUDA arrays的形式分配。CUDA array为了texure fetching做了优化，其内存结构不透明。Linear memory在device上有一个40-bit的地址空间，因此单独分配的实体可以通过指针访问彼此。

Linear memory典型的分配方式是通过cudaMalloc()来完成，然后由cudaFree()释放。通过cudaMemcpy()来完成host memory和device memory之间的数据转移。

除了上述函数之外，linear memory也可以通过cudaMallocPitch()和cudaMalloc3D()来分配，在分配2D或3D arrays时推荐使用这些函数，因为它们在分配过程中通过适当的padding来满足对齐要求（alignment requirements）。在访问array elements时必须使用返回的pitch或stride。

Shared Memory

通过__shared__限定符来分配shared memory，shared memory比global memory要快得多。

Multi-Device System

设备枚举

int deviceCount;
cudaGetDeviceCount(&deviceCount);
int device;
for (device = 0; device < deviceCount; ++device) {
    cudaDeviceProp deviceProp;
    cudaGetDeviceProperties(&deviceProp, device);
    printf("Device %d has compute capability %d.%d.\n",
            device, deviceProp.major, deviceProp.minor);
}

设备选择

Host thread随时可以通过cudaSetDevice()设置其要操作的device，device memory的分配和kernel的launche都在当前设置的device上进行。如果没有对cudaSetDevice()的调用，当前device为device 0。

Driver API

Driver API在cuda动态链接库（cuda.dll或cuda.so）中实现。所有的entry point均以cu开头。

Driver API是handle-based、命令式的API：大多数objects由opaque handles引用，这些handles可指定给操作这些objects的函数。

Table 15. Objects Available in the CUDA Driver API
Object Handle Description
Device CUdevice CUDA-enabled device
Context CUcontext Roughly equivalent to a CPU process
Module CUmodule Roughly equivalent to a dynamic library
Function CUfunction Kernel
Heap memory CUdeviceptr Pointer to device memory
CUDA array CUarray Opaque container for one-dimensional or two-dimensional data on the device, readable via texture or surface references
Texture reference CUtexref Object that describes how to interpret texture memory data
Surface reference CUsurfref Object that describes how to read or write CUDA arrays
Event CUevent Object that describes a CUDA event

Driver API在调用其中任何一个函数之前必须先用cuInit()初始化，然后必须创建一个关联到一个特定device上的CUDA context。

在一个CUDA context内，kernel被显式地以PTX或binary objects的形式加载。因此用C写成的kernel也必须单独编译成PTX或binary objects。

下面是一个使用driver API的示例：

int main() {
    int N = ...;
    size_t size = N * sizeof(float);

    // Allocate input vectors h_A and h_B in host memory
    float* h_A = (float*)malloc(size);
    float* h_B = (float*)malloc(size);

    // Initialize input vectors
    ...

    // Initialize
    cuInit(0);

    // Get number of devices supporting CUDA
    int deviceCount = 0;
    cuDeviceGetCount(&deviceCount);
    if (deviceCount == 0) {
        printf("There is no device supporting CUDA.\n");
        exit (0);
    }

    // Get handle for device 0
    CUdevice cuDevice;
    cuDeviceGet(&cuDevice, 0);

    // Create context
    CUcontext cuContext;
    cuCtxCreate(&cuContext, 0, cuDevice);

    // Create module from binary file
    CUmodule cuModule;
    cuModuleLoad(&cuModule, "VecAdd.ptx");

    // Allocate vectors in device memory
    CUdeviceptr d_A;
    cuMemAlloc(&d_A, size);
    CUdeviceptr d_B;
    cuMemAlloc(&d_B, size);
    CUdeviceptr d_C;
    cuMemAlloc(&d_C, size);

    // Copy vectors from host memory to device memory
    cuMemcpyHtoD(d_A, h_A, size);
    cuMemcpyHtoD(d_B, h_B, size);

    // Get function handle from module
    CUfunction vecAdd;
    cuModuleGetFunction(&vecAdd, cuModule, "VecAdd");

    // Invoke kernel
    int threadsPerBlock = 256;
    int blocksPerGrid =
        (N + threadsPerBlock - 1) / threadsPerBlock;
    void* args[] = { &d_A, &d_B, &d_C, &N };
    cuLaunchKernel(vecAdd, blocksPerGrid, 1, 1,
            threadsPerBlock, 1, 1, 0, 0, args, 0);
    ... 
}

Context

CUDA context类似于一个CPU process。driver API中所有资源以及执行的操作都被封装在一个CUDA context中，在一个context销毁时系统自动清理其中的资源。每个context拥有自己distinct地址空间，因此不同context中CUdeviceptr的值引用不同的内存位置。

一个host thread在任一时刻只能拥有一个device context。当通过cuCtxCreate()创建一个context时，该context被设置为调用该函数的host thread的current context。

Module

Modules是一些可动态加载的packages，包含device code和数据。

Kernel Execution

cuLaunchKernel()根据一个给定的execution configuration launch一个kernel。

参数传递有两种形式：

  • 通过一个指针数组传递，指针在数组中的位置对应着相应参数在参数列表中的位置，从指针指向的内存位置拷贝参数
  • 通过extra options

Others

cuMemcpyHtoD&cuMemcpyDtoH

CUresult cuMemcpyHtoD   (CUdeviceptr dstDevice, const void * srcHost, size_t ByteCount)
CUresult cuMemcpyDtoH   (void * dstHost, CUdeviceptr srcDevice, size_t ByteCount)
 CUresult 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
 )
  • f
    • Kernel to launch
  • gridDimX
    • Width of grid in blocks
  • gridDimY
    • Height of grid in blocks
  • gridDimZ
    • Depth of grid in blocks
  • blockDimX
    • X dimension of each thread block
  • blockDimY
    • Y dimension of each thread block
  • blockDimZ
    • Z dimension of each thread block
  • sharedMemBytes
    • Dynamic shared-memory size per thread block in bytes
  • hStream
    • Stream identifier
  • kernelParams
    • Array of pointers to kernel parameters
  • extra
    • Extra options
标签：