CUDA学习笔记(2) 第一个CUDA程序

  首先,我们用VisualStudio创建了CUDA的工程后,会出现一个数组对应位置元素求和的模版代码,我们可以先借此了解CUDA工程的主体结构,然后将他们全都删掉,从头开始练习。
  假设我们现在要创建384个线程,并要知道他们具体属于哪个线程束(Warp)线程块(Block),线程序号是多少。

  在 main() 函数中,我主要执行以下几个步骤:
1. 读取GPU硬件信息。
2. 计算。
3. 重置GPU。

  其中前两步需要我们自己具体设计,重置GPU可以直接调用官方的函数。

int main(int argc, char *argv[])
{
    cudaError_t cudaStatus;

    // 读取、检查设备信息
    check_Cuda_information(argc, argv);

    // 计算部分
    cudaStatus = caculate_Cuda_function();

    // 重置设备以便 Nsight 、 Visual Profiler 记录运行时间
    cudaStatus = cudaDeviceReset();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceReset failed!");
        return 1;
    }

    return 0;
}


  第一步,我们需要包含对应的头文件,并添加对应的包含目录。

// CUDA
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
// includes CUDA Runtime
#include <cuda_runtime.h>
// includes, project
#include <helper_cuda.h>
#include <helper_functions.h>
// C IO
#include <stdio.h>
// C++ IOstream
#include <iostream>
using namespace std;


  第二步,我们读取GPU的硬件属性,并记录每个线程块的最大线程数

// 块最大线程数
int max_thread_per_block = 0;
// CUDA检查设备信息
void check_Cuda_information(int main_argc, char ** main_argv);

// 检查显卡硬件属性
void check_Cuda_information(int main_argc, char ** main_argv)
{
    // 设备ID
    int devID;
    // 设备属性
    cudaDeviceProp deviceProps;

    // 获取设备ID
    devID = findCudaDevice(main_argc, (const char **)main_argv);

    // 获取GPU信息
    checkCudaErrors(cudaGetDeviceProperties(&deviceProps, devID));
    cout << "devID = " << devID << endl;
    // 显卡名称
    cout << "CUDA device is \t\t\t" << deviceProps.name << endl;
    // 每个 线程块(Block)中的最大线程数
    cout << "CUDA max Thread per Block is \t" << deviceProps.maxThreadsPerBlock << endl;
    max_thread_per_block = deviceProps.maxThreadsPerBlock;
}


  第三步,编写一个CUDA计算函数,执行以下步骤:

1.定义主机变量(Host,指CPU部分及内存中的数据),定义设备(Device,指GPU及显存中的数据)变量。
cudaError_t cudaStatus; 用来接收CUDA官方函数的返回值,检验函数是否正确执行。

// CUDA计算部分
// CUDA计算部分
// Helper function for using CUDA to add vectors in parallel.
cudaError_t caculate_Cuda_function()
{
    cudaError_t cudaStatus;
    // my_check_CUDA_status 调用计数,方便调试查错
    int use_counter = 0;
    // Host变量(内存)
    int ARRAY_LENGTH = 3 * 2 * 64;
    int *thread_index, *warp_index, *block_index;

    thread_index = (int*)malloc(ARRAY_LENGTH * sizeof(int));
    warp_index = (int*)malloc(ARRAY_LENGTH * sizeof(int));
    block_index = (int*)malloc(ARRAY_LENGTH * sizeof(int));

    // Device变量(显存)
    int *dev_thread_index = NULL, *dev_warp_index = NULL, *dev_block_index = NULL;


2.选择我们准备使用的GPU。

    // 选择我们准备使用的设备,在有多块GPU的电脑中,这一步十分重要!
    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }


3.分配显存,cudaMalloc有两个输入参数,一个返回值。

  • 第一个是 指针变量的地址 ,指针变量dev_thread_index本身存放在内存中,而他的内容是显存的地址!
  • 第二个是分配的显存大小。
  • 返回值表示函数是否正确的执行。
    // 分配显存
    cudaStatus = cudaMalloc((void**)&dev_thread_index, ARRAY_LENGTH * sizeof(int));
    if (!my_check_CUDA_status(cudaStatus, use_counter)) goto Error;
    cudaStatus = cudaMalloc((void**)&dev_warp_index, ARRAY_LENGTH * sizeof(int));
    if (!my_check_CUDA_status(cudaStatus, use_counter)) goto Error;
    cudaStatus = cudaMalloc((void**)&dev_block_index, ARRAY_LENGTH * sizeof(int));
    if (!my_check_CUDA_status(cudaStatus, use_counter)) goto Error;

  这里我定义一个内联函数 my_check_CUDA_status 来判断GPU是否正确运行,如果发生错误,则输出在第几次调用这个函数时发生的错误,以便快速定位错误的代码。

// 检查指令在GPU是否正确运行
inline bool my_check_CUDA_status(cudaError_t inline_cudaStatus, int & use_counter)
{
    use_counter++;

    if (inline_cudaStatus != cudaSuccess)
    {
        fprintf(stderr, "CUDA failed! use_counter = %d\r\n", use_counter);
        cout << "inline_cudaStatus = " << inline_cudaStatus << endl;
        return false;
    }
    else
    {
        return true;
    }
}


4.在编写核函数时我们一般会用到几个CUDA已经定义好的变量:

变量名说明
threadIdx.x, threadIdx.y, threadIdx.z线程(Thread) x、y、z三个维度的下标
blockIdx.x, blockIdx.y, blockIdx.z线程块(Block) x、y、z三个维度的下标
blockDim.x, blockDim.y, blockDim.z一个线程块(Block)单元中x、y、z三个维度的线程(Thread)的数量
gridDim.x, gridDim.y, gridDim.z一个线程网格(Grid)单元中x、y、z三个维度的线程块(Block)的数量
warpSize线程束(Warp)的大小(一般为32)


  这里举一个线程块(Block)线程网格(Grid)都是二维情况的例子:
这里写图片描述

  在核函数中__global__前缀表示这个函数或变量是一个全局的,int block_indexint thread_index作为每个流处理器的寄存器变量,分别记录当前线程的线程序号与所在的线程块序号。核函数必须为void类型且不能有return!

__global__ void Kernel_func(int * thread_index_array, int * warp_index_array, int * block_index_array)
{
    int block_index = blockIdx.x + blockIdx.y * gridDim.x + blockIdx.z * gridDim.x * gridDim.y;
    int thread_index = block_index * blockDim.x * blockDim.y * blockDim.z + \
        threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y;

    thread_index_array[thread_index] = thread_index;
    warp_index_array[thread_index] = thread_index / warpSize;
    block_index_array[thread_index] = block_index;
}

  那么如何编写一个合理、高效的核函数呢,我们必须了解以下几个GPU底层的工作原理。


这里写图片描述

  在GPU中,每32个线程作为一个线程束(Warp)一起执行,一个线程束(Warp)同时执行读取显存等操作,所以线程块(Block)中一个单元的大小最好大于32且为32的整数倍,这样才能尽可能让GPU高效地运行。
  如上图所示(参考自1),128个线程分为4个线程束,线程束0到3依次被处理内存请求,只需要4次周期操作就可以完成对显存数据的读写。若一个线程块中只有一个线程,那么就会有128次内存请求,这样使得程序的执行大打折扣。

  NVIDIA公司在编写CUDA的时候为用户提供了一个叫做 dim3 的结构体,用来定义线程块(Block)线程网格(Grid)每一个单元的大小。下面是源码,我们可以看到它其实是由3个int组成的结构体,分别代表x、y、z三个维度的下标最大值。

struct __device_builtin__ dim3
{
    unsigned int x, y, z;
#if defined(__cplusplus)
    __host__ __device__ dim3(unsigned int vx = 1, unsigned int vy = 1, unsigned int vz = 1) : x(vx), y(vy), z(vz) {}
    __host__ __device__ dim3(uint3 v) : x(v.x), y(v.y), z(v.z) {}
    __host__ __device__ operator uint3(void) { uint3 t; t.x = x; t.y = y; t.z = z; return t; }
#endif /* __cplusplus */
};

  CUDA的核函数的调用方法如下,规定使用“<<<”和“>>>”符号来写CUDA必要的参数列表,即线程块(Block)线程网格(Grid)每一个单元的大小。后面圆括号中的内容是我们需要传入的显存内数据的地址(指针)。

    // 定义网格的大小(block_rect)、块的大小(thread_rect)
    dim3 block_rect(3, 2, 1), thread_rect(64, 1, 1);
    cout << "block_rect :\t" << block_rect.x << "\t" << block_rect.y << "\t" << block_rect.z << "\t" << endl;
    cout << "thread_rect :\t" << thread_rect.x << "\t" << thread_rect.y << "\t" << thread_rect.z << "\t" << endl;

    // GPU开始计算(传入核函数)
    Kernel_func <<< block_rect, thread_rect >>>(dev_thread_index, dev_warp_index, dev_block_index);
  • block_rect是指线程网格(Grid)的一个单元中容纳block各维度的数量。
  • thread_rect是线程块(Block)一个单元中容纳thread各维度的数量。当然线程块(Block)一个单元中容纳thread各维度的数量不能超过对应架构GPU所规定的量,即maxThreadsPerBlock。
      例如maxThreadsPerBlock为1024,那么 (thread_rect.x * thread_rect.y * thread_rect.z)的值不能超过1024。

这里写图片描述

5.检查核函数是否正常运行。

    // 检查核函数执行是否报错
    cudaStatus = cudaGetLastError();
    if (!my_check_CUDA_status(cudaStatus, use_counter)) goto Error;

    // 与GPU同步,并检查是否出现错误
    cudaStatus = cudaDeviceSynchronize();
    if (!my_check_CUDA_status(cudaStatus, use_counter)) goto Error;


6.传出数据

// 传出数据
    cudaStatus = cudaMemcpy(thread_index, dev_thread_index, ARRAY_LENGTH * sizeof(int), cudaMemcpyDeviceToHost);
    if (!my_check_CUDA_status(cudaStatus, use_counter)) goto Error;
    cudaStatus = cudaMemcpy(warp_index, dev_warp_index, ARRAY_LENGTH * sizeof(int), cudaMemcpyDeviceToHost);
    if (!my_check_CUDA_status(cudaStatus, use_counter)) goto Error;
    cudaStatus = cudaMemcpy(block_index, dev_block_index, ARRAY_LENGTH * sizeof(int), cudaMemcpyDeviceToHost);
    if (!my_check_CUDA_status(cudaStatus, use_counter)) goto Error;

7.错误处理及内存、显存释放。

Error:
    // 释放显存
    cudaFree(dev_thread_index);
    cudaFree(dev_warp_index);
    cudaFree(dev_block_index);
    // 释放内存
    free(thread_index);
    free(block_index);
    free(warp_index);

总结:

  我认为使用CUDA的过程可以分为以下几个步骤:
1. 读取GPU硬件信息,以保证我们的代码可以兼容、高效地在不同型号GPU上运行。
2. 选择我们要使用的GPU序号,这一步在有多块GPU的平台上尤为重要。
3. 申请全局显存。
4. 调用一次核函数。
5. CPU与GPU同步(通信),以检查GPU计算过程是否出错,并读出所需数据。
6. 释放显存,重置GPU。

  值得注意的是,每调用一次核函数,就必须执行一次步骤5检查执行状况,之后才能再调用一次核函数。





我的 learn_CUDA_02.cu 完整代码:

// CUDA
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
// includes CUDA Runtime
#include <cuda_runtime.h>
// includes, project
#include <helper_cuda.h>
#include <helper_functions.h> // helper utility functions 

// C IO
#include <stdio.h>
// C++ IOstream
#include <iostream>
using namespace std;


// CUDA检查设备信息
void check_Cuda_information(int main_argc, char ** main_argv);
// CUDA计算部分
cudaError_t caculate_Cuda_function();
// 块最大线程数
int max_thread_per_block = 0;

__global__ void Kernel_func(int * thread_index_array, int * warp_index_array, int * block_index_array)
{
    int block_index = blockIdx.x + blockIdx.y * gridDim.x + blockIdx.z * gridDim.x * gridDim.y;
    int thread_index = block_index * blockDim.x * blockDim.y * blockDim.z + \
        threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y;

    thread_index_array[thread_index] = thread_index;
    warp_index_array[thread_index] = thread_index / warpSize;
    block_index_array[thread_index] = block_index;
}


int main(int argc, char *argv[])
{
    cudaError_t cudaStatus;

    // 读取、检查设备信息
    check_Cuda_information(argc, argv);

    // 计算部分
    cudaStatus = caculate_Cuda_function();

    // cudaDeviceReset must be called before exiting in order for profiling and
    // tracing tools such as Nsight and Visual Profiler to show complete traces.
    cudaStatus = cudaDeviceReset();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceReset failed!");
        return 1;
    }

    return 0;
}


// 检查显卡硬件属性
void check_Cuda_information(int main_argc, char ** main_argv)
{
    // 设备ID
    int devID;
    // 设备属性
    cudaDeviceProp deviceProps;

    //
    cout << "argc = " << main_argc << endl;
    for (int i = 0; i < main_argc; i++)
    {
        printf("argv[%d] = %s\r\n", i, main_argv[i]);
    }
    cout << endl;

    // 获取设备ID
    devID = findCudaDevice(main_argc, (const char **)main_argv);

    // 获取GPU信息
    checkCudaErrors(cudaGetDeviceProperties(&deviceProps, devID));
    cout << "devID = " << devID << endl;
    // 显卡名称
    cout << "CUDA device is \t\t\t" << deviceProps.name << endl;
    // 每个 线程块(Block)中的最大线程数
    cout << "CUDA max Thread per Block is \t" << deviceProps.maxThreadsPerBlock << endl;
    max_thread_per_block = deviceProps.maxThreadsPerBlock;
    // 每个 多处理器组(MultiProcessor)中的最大线程数
    cout << "CUDA max Thread per SM is \t" << deviceProps.maxThreadsPerMultiProcessor << endl;
    // GPU 中 SM 的数量
    cout << "CUDA SM counter\t\t\t" << deviceProps.multiProcessorCount << endl;
    // 线程束大小
    cout << "CUDA Warp size is \t\t" << deviceProps.warpSize << endl;
    // 每个SM中共享内存的大小
    cout << "CUDA shared memorize is \t" << deviceProps.sharedMemPerMultiprocessor << "\tbyte" << endl;
}

// 检查指令在GPU是否正确运行
inline bool my_check_CUDA_status(cudaError_t inline_cudaStatus, int & use_counter)
{
    use_counter++;

    if (inline_cudaStatus != cudaSuccess)
    {
        fprintf(stderr, "CUDA failed! use_counter = %d\r\n", use_counter);
        cout << "inline_cudaStatus = " << inline_cudaStatus << endl;
        return false;
    }
    else
    {
        return true;
    }
}

// CUDA计算部分
// Helper function for using CUDA to add vectors in parallel.
cudaError_t caculate_Cuda_function()
{
    cudaError_t cudaStatus;
    // my_check_CUDA_status 调用计数,方便调试查错
    int use_counter = 0;
    // Host变量(内存)
    const int ARRAY_LENGTH = 3 * 2 * 64;
    int *thread_index, *warp_index, *block_index;

    thread_index = (int*)malloc(ARRAY_LENGTH * sizeof(int));
    warp_index = (int*)malloc(ARRAY_LENGTH * sizeof(int));
    block_index = (int*)malloc(ARRAY_LENGTH * sizeof(int));

    // Device变量(显存)
    int *dev_thread_index = NULL, *dev_warp_index = NULL, *dev_block_index = NULL;


#pragma region(选择GPU)
    // Choose which GPU to run on, change this on a multi-GPU system.
    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }
#pragma endregion


#pragma region(分配显存、传入数据)
    // 分配显存
    cudaStatus = cudaMalloc((void**)&dev_thread_index, ARRAY_LENGTH * sizeof(int));
    if (!my_check_CUDA_status(cudaStatus, use_counter)) goto Error;
    cudaStatus = cudaMalloc((void**)&dev_warp_index, ARRAY_LENGTH * sizeof(int));
    if (!my_check_CUDA_status(cudaStatus, use_counter)) goto Error;
    cudaStatus = cudaMalloc((void**)&dev_block_index, ARRAY_LENGTH * sizeof(int));
    if (!my_check_CUDA_status(cudaStatus, use_counter)) goto Error;

#pragma endregion


#pragma region(执行核函数,并在核函数完成时检查错误报告)
    // 定义网格的大小(block_rect)、块的大小(thread_rect)
    dim3 block_rect(3, 2, 1), thread_rect(64, 1, 1);
    cout << "block_rect :\t" << block_rect.x << "\t" << block_rect.y << "\t" << block_rect.z << "\t" << endl;
    cout << "thread_rect :\t" << thread_rect.x << "\t" << thread_rect.y << "\t" << thread_rect.z << "\t" << endl;

    // GPU开始计算(传入核函数)
    Kernel_func <<< block_rect, thread_rect >>>(dev_thread_index, dev_warp_index, dev_block_index);

    // 检查核函数执行是否报错
    cudaStatus = cudaGetLastError();
    if (!my_check_CUDA_status(cudaStatus, use_counter)) goto Error;

    // 与GPU同步,并检查是否出现错误
    cudaStatus = cudaDeviceSynchronize();
    if (!my_check_CUDA_status(cudaStatus, use_counter)) goto Error;

#pragma endregion


#pragma region(传出数据)
    // 传出数据
    cudaStatus = cudaMemcpy(thread_index, dev_thread_index, ARRAY_LENGTH * sizeof(int), cudaMemcpyDeviceToHost);
    if (!my_check_CUDA_status(cudaStatus, use_counter)) goto Error;
    cudaStatus = cudaMemcpy(warp_index, dev_warp_index, ARRAY_LENGTH * sizeof(int), cudaMemcpyDeviceToHost);
    if (!my_check_CUDA_status(cudaStatus, use_counter)) goto Error;
    cudaStatus = cudaMemcpy(block_index, dev_block_index, ARRAY_LENGTH * sizeof(int), cudaMemcpyDeviceToHost);
    if (!my_check_CUDA_status(cudaStatus, use_counter)) goto Error;

#pragma endregion


    for (int i = 0; i < ARRAY_LENGTH; i++)
    {
        printf("thread index \t: %d\t", thread_index[i]);
        printf("warp flag \t: %d\t", warp_index[i]);
        printf("block index \t: %d\t\r\n", block_index[i]);
    }


Error:
    // 释放显存
    cudaFree(dev_thread_index);
    cudaFree(dev_warp_index);
    cudaFree(dev_block_index);
    // 释放内存
    free(thread_index);
    free(block_index);
    free(warp_index);


    return cudaStatus;
}

运行结果:

这里写图片描述





参考:

1.《CUDA并行程序设计》机械工业出版社

2.Fermi架构白皮书

3.CUDA C Programming Guide

参与评论 您还未登录,请先 登录 后发表或查看评论

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
©️2022 CSDN 皮肤主题:编程工作室 设计师:CSDN官方博客 返回首页

打赏作者

Bingjian-Gong

你的鼓励将是我创作的最大动力

¥2 ¥4 ¥6 ¥10 ¥20
输入1-500的整数
余额支付 (余额:-- )
扫码支付
扫码支付:¥2
获取中
扫码支付

您的余额不足,请更换扫码支付或充值

打赏作者

实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、C币套餐、付费专栏及课程。

余额充值