host代码与device代码

  • host代码(主机端)指运行在CPU及内存(RAM)上的代码，一般是标准C/CPP代码

  • device代码(设备端)指运行在GPU及显存(vRAM)上的代码，即声明为__global__的代码(kernel函数)

一个cuda c程序由host代码和device代码两部分组成，使用nvcc编译一个.cu文件时，nvcc将host代码交给host compiler(比如msvc)当做正常的C/CPP代码编译，将device代码交给cuda编译器编译。

示例如下：

#include <iostream>

__global__ void add(int a, int b, int *c){
    *c = a + b;
}

int main(){
    int c; int *dev_c;

    cudaMalloc((void**)&dev_c, sizeof(int));

    add<<<1,1>>>(2,7,dev_c);

    cudaMemcpy(&c,dev_c,sizeof(int),cudaMemcpyDeviceToHost);

    std::cout << c << std::endl;

    return 0;
}

我们逐行来看：

  1. 被__global__标识的函数就是一个kernel函数，即device代码，它实际上是指派给GPU的异步任务，虽然使用上很像，但也不能完全当平时的C/CPP函数来看待。一个很明显的区别是：kernel函数必须设置为void，因为host和device没有共同的调用栈帧，kernel函数无法直接给host端返回数据。

  2. cudaMalloc()是cuda runtime提供的接口，和malloc()类似，它可以分配显存。cudaMalloc((void**)&dev_c, sizeof(int)) 分配了一个int大小的显存，dev_c 就成为了指向这个显存的device指针。device指针只能由cudaMalloc()分配，由cudaFree()释放，由cudaMemcpy()复制内容。不能直接在host端解引用device指针。

  3. 调用kernel函数时，需要以类似<<<1,1>>>的方式传递一组参数。它不是传给device code的，是告诉cuda编译器如何执行device code的。这些参数如何设置会在下章学。

由于kernel函数没有返回值，host端也不能直接对device指针解引用，因此host端与device端的通信一般通过cudaMemcpy()或后面要提到的其他手段。

从上面的代码可以看出来，cuda代码实际上就是在cpp中内嵌了device代码。这么自然地对device代码调用是cuda编译器的功劳，背后的原理很复杂，但展现在我们面前的开发方式却很简单。

查询和使用device信息

查询device信息主要用到以下两个api：

cudaGetDeviceCount(int *count); //查询cuda设备数量
cudaGetDeviceProperties(cudaDeviceProp *prop, int device); //查询cuda设备信息

其中prop是一个结构体，可以用如下方式使用：

#include <stdio.h>

int main(void) {
    int count;
    cudaGetDeviceCount(&count)
    printf("共有 %d 个 CUDA 设备\n", count);

    for (int i = 0; i < count; i++) {
        cudaDeviceProp prop;
        cudaGetDeviceProperties(&prop, i)
        printf("设备 #%d: %s\n", i, prop.name);
        printf("  全局内存大小: %zu 字节\n", prop.totalGlobalMem);
        printf("  每个 block 最大线程数: %d\n", prop.maxThreadsPerBlock);
        printf("  多处理器数量: %d\n", prop.multiProcessorCount);
        printf("  Warp 大小: %d\n", prop.warpSize);
        printf("  时钟频率: %.2f MHz\n", prop.clockRate / 1000.0f);
        printf("\n");
    }
}

通过prop中的各种信息，可以应对多显卡，以及编写可移植的cuda程序等情况。(根据显卡prop不同做不同的处理)