3维线程格 gpu_CUDA编程3-编程模型（下）

3 组织并行线程

使用合适的网格和块大小来正确地组织线程，对内核会产生很大的影响。本节通过一个矩阵加法的例子，来验证这一点。对于矩阵运算，传统的方法是使用一个包含二维网格与二维块的布局来组织线程，但这并样性能并不一定最佳，下面探索几种线程布局方法，来探索网格和块的启发性用法。

3.1 建立矩阵索引

通常情况一个矩阵用行优先方法在全局内存中进行线性存储，如下图的例子。在一个矩阵加法核函数中，一个线程通常被分配一个数据元素来处理，首先要完成的是从块和线程索引中得到数据在全局内存中的位置。对一个二维示例来说通常有需要3种索引：

• 线程和块索引；
• 矩阵给定点的坐标；
• 全局线性内存中的偏移量。

对于一个给定线程，首先可以将线程和块索引映射到矩阵坐标，以获取线程块和线程索引的全局内存偏移量；然后将这些矩阵坐标映射到全局内存的存储单元中。具体方法如下：

1. 首先，用以下公式把线程和块索引映射到矩阵坐标上

ix = threadIdx.x + blockIdx.x * blockDim.x;
iy = threadIdx.y + blockIdx.y * blockDim.y;

2. 然后使用下面的公式把矩阵坐标映射到全局内存索引中：

idx = iy * nx + ix;

原理如下图所示：

3.2 不同布局的矩阵求和

下面addMatrix.cu的代码，分别使用了二维网格+二维块，一维网格+一维块，以及二维网格+一维块三种线程组织的形式。 通过在第一个命令行参数输入0、1、2可以不分别做选择这种组织形式。

#include <stdio.h>

void addMatrixOnHost(float *A, float *B, float *C, const int nx, const int ny)
{
    float *ia = A;
    float *ib = B;
    float *ic = C;
    for(int iy = 0; iy < ny; ++iy)
    {
        for(int ix = 0; ix < nx; ++ix)
            ic[ix] = ia[ix] + ib[ix];
        ia += nx;
        ib += nx;
        ic += nx;
    }
}

__global__ void addMatrixOnGPU2D(float *A, float *B, float *C, int nx, int ny)
{
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int idx = iy * nx + ix;
    if(ix < nx && iy < ny)
        C[idx] = A[idx] + B[idx];
}

__global__ void addMatrixOnGPU1D(float *A, float *B, float *C, int nx, int ny)
{
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
    if(ix < nx)
        for(int iy = 0; iy < ny; ++iy)
        {
            int idx = iy * nx + ix;
            C[idx] = A[idx] + B[idx];
        }
}

__global__ void addMatrixOnGPUMD(float *A, float *B, float *C, int nx, int ny)
{
    unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int iy = threadIdx.y;
    unsigned int idx = iy * nx + ix;
    if(ix < nx && iy < ny)
        C[idx] = A[idx] + B[idx];
}

int main(int argc, char **argv)
{
    printf("%s Starting...n", argv[0]);

    // 设置设备
    int dev = 0;
    cudaDeviceProp devProp;
    CHECK(cudaGetDeviceProperties(&devProp, dev));
    printf("Using Device %d: %sn", dev, devProp.name);
    CHECK(cudaSetDevice(dev));

    // 设置矩阵大小
    int nx = 1 << 14;
    int ny = 1 << 14;
    int nxy = nx * ny;
    int nBytes = nxy * sizeof(float);
    printf("Matrix size: %dx%d", nx, ny);

    // 设置主机内存
    float *h_A, *h_B, *hostRef, *gpuRef;
    h_A     = (float*)malloc(nBytes);
    h_B     = (float*)malloc(nBytes);
    hostRef = (float*)malloc(nBytes);
    gpuRef  = (float*)malloc(nBytes);
    initData(h_A, nxy);
    initData(h_B, nxy);
    memset(hostRef, 0, nBytes);
    memset(gpuRef, 0, nBytes);
    addMatrixOnHost(h_A, h_B, hostRef, nx, ny);

    // 设置设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc((float**)&d_A, nBytes);
    cudaMalloc((float**)&d_B, nBytes);
    cudaMalloc((float**)&d_C, nBytes);
    cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice);

    // 运行核函数
    double iStart, iElaps;
    int choice = atoi(argv[1]);
    // 二维x二维线程组织形式
    if(choice == 0)
    {
        int dimx = atoi(argv[2]);
        int dimy = atoi(argv[3]);
        dim3 block(dimx, dimy);
        dim3 grid((nx+block.x-1)/block.x, (ny+block.y-1)/block.y);
        iStart = cpuSecond();
        addMatrixOnGPU2D<<<grid, block>>>(d_A, d_B, d_C, nx, ny);
        cudaDeviceSynchronize();
        iElaps = cpuSecond() - iStart;
        printf("Execute the 2Dx2D layout of threads.n");
        printf("addMatrixOnGPU <<<(%d,%d), (%d,%d)>>> Time elapsed %f secondsn",
               grid.x, grid.y, block.x, block.y, iElaps);
    }
    // 一维x一维线程组织形式
    else if(choice == 1)
    {
        int dim = atoi(argv[2]);
        dim3 block(dim);
        dim3 grid((nx+block.x-1)/block.x);
        iStart = cpuSecond();
        addMatrixOnGPU1D<<<grid, block>>>(d_A, d_B, d_C, nx, ny);
        cudaDeviceSynchronize();
        iElaps = cpuSecond() - iStart;
        printf("Execute the 1Dx1D layout of threads.n");
        printf("addMatrixOnGPU <<<%d, %d>>> Time elapsed %f secondsn",
               grid.x, block.x, iElaps);
    }
    // 二维x一维线程组织形式
    else
    {
        int dim = atoi(argv[2]);
        dim3 block(dim);
        dim3 grid((nx+block.x-1)/block.x, ny);
        iStart = cpuSecond();
        addMatrixOnGPUMD<<<grid, block>>>(d_A, d_B, d_C, nx, ny);
        cudaDeviceSynchronize();
        iElaps = cpuSecond() - iStart;
        printf("Execute the 1Dx1D layout of threads.n");
        printf("addMatrixOnGPU <<<(%d,%d), %d>>> Time elapsed %f secondsn",
               grid.x, grid.y, block.x, iElaps);
    }

    // 验证结果
    cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);
    checkResult(hostRef, gpuRef, nxy);
    // 释放设备内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    // 释放主机内存
    free(h_A);
    free(h_B);
    free(hostRef);
    free(gpuRef);

    return 0;
}

从这个例子可以看出：

• 改变执行配置对内核性能有影响；
• 传统的核函数实现一般不能获得最佳性能；
• 对给定核函数尝试使用不同的网格和线程块大小可以获得更好的性能。

4 设备管理

NVIDIA提供了几种查询和管理GPU设备的方法，学会如何查询GPU设备，可以在运行时使用它来帮助设置内核的执行配置。本节介绍两种设备查询方法，一种是CUDA Runtime API，另一种是系统管理界面命令行程序nvidia-smi。

4.1 使用API查询信息

使用cudaGetDeviceProperties函数可以查询关于GPU设备的所有信息，其原型如下：

cudaError_t cudaGetDeviceProperties(cudaDeviceProp* prop, int device);

cudaDeviceProp结构体返回GPU设备的属性，下面的checkDeviceInfo.cu代码提供了一个示例：

#include <cuda_runtime.h>
#include <stdio.h>

int main(int argc, char **argv)
{
    printf("%s Starting...n", argv[0]);

    int deviceCount = 0;
    cudaGetDeviceCount(&deviceCount);
    if (deviceCount == 0)
        printf("There are no available device(s) that support CUDAn");
    else
        printf("Detected %d CUDA Capable device(s)n", deviceCount);

    int dev = 0, driverVersion = 0, runtimeVersion = 0;
    CHECK(cudaSetDevice(dev));
    cudaDeviceProp deviceProp;
    CHECK(cudaGetDeviceProperties(&deviceProp, dev));
    printf("Device %d: "%s"n", dev, deviceProp.name);

    cudaDriverGetVersion(&driverVersion);
    cudaRuntimeGetVersion(&runtimeVersion);
    printf("  CUDA Driver Version / Runtime Version          %d.%d / %d.%dn",
           driverVersion / 1000, (driverVersion % 100) / 10,
           runtimeVersion / 1000, (runtimeVersion % 100) / 10);
    printf("  CUDA Capability Major/Minor version number:    %d.%dn",
           deviceProp.major, deviceProp.minor);
    printf("  Total amount of global memory:                 %.2f GBytes (%llu "
           "bytes)n", (float)deviceProp.totalGlobalMem / pow(1024.0, 3),
           (unsigned long long)deviceProp.totalGlobalMem);
    printf("  GPU Clock rate:                                %.0f MHz (%0.2f "
           "GHz)n", deviceProp.clockRate * 1e-3f,
           deviceProp.clockRate * 1e-6f);
    printf("  Memory Clock rate:                             %.0f Mhzn",
           deviceProp.memoryClockRate * 1e-3f);
    printf("  Memory Bus Width:                              %d-bitn",
           deviceProp.memoryBusWidth);

    if (deviceProp.l2CacheSize)
    {
        printf("  L2 Cache Size:                                 %d bytesn",
               deviceProp.l2CacheSize);
    }

    printf("  Max Texture Dimension Size (x,y,z)             1D=(%d), "
           "2D=(%d,%d), 3D=(%d,%d,%d)n", deviceProp.maxTexture1D,
           deviceProp.maxTexture2D[0], deviceProp.maxTexture2D[1],
           deviceProp.maxTexture3D[0], deviceProp.maxTexture3D[1],
           deviceProp.maxTexture3D[2]);
    printf("  Max Layered Texture Size (dim) x layers        1D=(%d) x %d, "
           "2D=(%d,%d) x %dn", deviceProp.maxTexture1DLayered[0],
           deviceProp.maxTexture1DLayered[1], deviceProp.maxTexture2DLayered[0],
           deviceProp.maxTexture2DLayered[1],
           deviceProp.maxTexture2DLayered[2]);
    printf("  Total amount of constant memory:               %lu bytesn",
           deviceProp.totalConstMem);
    printf("  Total amount of shared memory per block:       %lu bytesn",
           deviceProp.sharedMemPerBlock);
    printf("  Total number of registers available per block: %dn",
           deviceProp.regsPerBlock);
    printf("  Warp size:                                     %dn",
           deviceProp.warpSize);
    printf("  Maximum number of threads per multiprocessor:  %dn",
           deviceProp.maxThreadsPerMultiProcessor);
    printf("  Maximum number of threads per block:           %dn",
           deviceProp.maxThreadsPerBlock);
    printf("  Maximum sizes of each dimension of a block:    %d x %d x %dn",
           deviceProp.maxThreadsDim[0],
           deviceProp.maxThreadsDim[1],
           deviceProp.maxThreadsDim[2]);
    printf("  Maximum sizes of each dimension of a grid:     %d x %d x %dn",
           deviceProp.maxGridSize[0],
           deviceProp.maxGridSize[1],
           deviceProp.maxGridSize[2]);
    printf("  Maximum memory pitch:                          %lu bytesn",
           deviceProp.memPitch);

    exit(EXIT_SUCCESS);
}

编译运行：

nvcc checkDeviceInfo.cu -o checkdev
./checkdev

在本人实验的机器上运行结果为：

4.2 使用nvidia-smi查询GPU信息

nvidia-smi是一个命令行工具，用于管理和监控GPU设备，并允许查询和修改设备状态。比如，要查询系统中安装了多少个GPU设备以及每个设备的ID，可以使用如下命令：

nvidia-smi -L

可以使用下面的命令获取GPU 0的详细信息

nvidia-smi -q -i 0

可以使用下面的参数精简nvidia-smi的显式信息：

• MEMORY
• UTILIZATION
• ECC
• TEMPERATURE
• POWER
• CLOCK
• COMPUTE
• PIDS
• PERFORMANCE
• SUPPORTED_CLOCKS
• PAGE_RETIREMENT
• ACCOUNTING

比如，只显示设备内存信息，可以使用：

nvidia-smi -q -i 0 -d MEMORY | tail -n 5

只显示设备使用信息，可以使用：

nvidia-smi -q -i 0 -d UTILIZATION | tail -n 4

4.3 在运行时设置设备

在一个有N个GPU的系统中，可以使用环境变量CUDA_VISIBLE_DEVICES 来使用指定ID的设备，而无需更改程序。

如果运行时设置CUDA_VISIBLE_DEVICES=2，nvidia驱动程序会屏蔽其他GPU，这是设备2作为设备0出现在程序中。

也可以指定多个设备，比如设置CUDA_VISIBLE_DEVICES =2,3，运行时Nvidia驱动程序将只使用ID为2和3的设备，且分别映射为0和1。