CUDA编程实战-1

最新推荐文章于 2025-06-25 10:29:33 发布

Riesman

最新推荐文章于 2025-06-25 10:29:33 发布

阅读量1.2k

点赞数

CC 4.0 BY-SA版权

分类专栏： c语言文章标签： cuda

本文链接：https://blog.youkuaiyun.com/u010458863/article/details/59106655

c语言专栏收录该内容

13 篇文章

订阅专栏

前言：
GPU作为如今主流的异构计算平台，收到越来越多的追捧，尤其是深度学习的发展，更加催生了其应用。作为并行计算领域的一枚，必须要好好学习一番啦。

GPU各种型号整理

G80、 Fermi、Kepler、 Maxwell 和 Pascal 5 代架构整理：TODO

安装

TODO

基础知识

（1）函数的类型。
__host__ float HostFunc():默认情况。被host函数调用，在host函数上执行。（host指cpu端）
__devide__ float DeviceFunc(); 被设备函数调用，在设备上执行。（设备指gpu）
__global__ void KernelFunc(): 被host函数调用，在设备上执行。
注意:
（a）__global__ 函数返回值必须为void
（b）设备上执行的函数：不能是递归的，函数参数必须固定，不能再函数内部使用static变量。

（2）变量类型
__shared__ A[4]；//在share memory，块内线程共享。
设备上的函数，声明的变量都是存在register上的，存不下的放到local memory；
cudaMalloc()的空间是在设备的global memory上的。
(3)几个cuda头文件

    #include<cuda_runtime.h>:  里面包含cuda 的基本routine
    #include <helper_cuda.h>：里面包含error处理等routine
    #include <helper_functions.h>： 这里面是什么？？？

几个例子

CUDA安装完后，默认安装路径（下文成为：$CUDA_ROOT）是：

/usr/local/cuda-8.0（假设安装的是8.0）

在$CUDA_ROOT下的samples下，有很多例子。先从这些例子学起。

学习列表

1. vectorAdd.cu:实现向量加（一位dim） 

2. cudaOpenMP.cu：使用openmp多线程，调用不同的gpu（多gpu编程） 

3. simpleMPI：mpi与cuda的结合。跟孙乔跟我说的差不多。mpi负责分配任务，然后每个任务都是在cuda上执行的。设计的很分离，不耦合。good

4. simplePrintf：在cuda上打印。

5. simpleAssert：通过是第几个线程，若线程数超过N， 就assert error

6. matrixMul：分块的矩阵乘，用到了share memory。

7. matrixMulCUBLAS：调用cublas计算gemm的例子。

8. samples/6_Advanced/scan：good。看看这个scan是如何实现的

9. samples/6_Advanced/ reduction：

10. samples/6_Advanced/ transpose ：

11. samples/6_Advanced/ mergeSort：

12. samples/3_Imaging/convolutionFFT2D : fft版本的卷积

例子讲解

实战1只介绍前5个函数。我们把cuda routine挑出来。挨个学习。
变量：

    1. cudaError_t err = cudaSuccess;
    //cudaError_t类型，表示错误类型。cudaSuccess表示成功。一般cuda routine的返回值都是cudaError_t类型，表示函数是否执行成功。  

    2. printf("%s\n", cudaGetErrorString(cudaGetLastError()));
    //输出错误时，使用以上函数转化为string。

函数：
1. err = cudaMalloc((void **)&d_A, size);
//动态内存申请函数，在设备的global memory上申请size个字节空间。

2. err = cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);or
    err = cudaMemcpy(h_A, d_A, size, cudaMemcpyDeviceToHost);
    //内存拷贝函数：从cpu上的内存h_A上拷贝size个字节数据到gpu上的内存d_A。反之，一样。

3. int threadsPerBlock = 256;
    int blocksPerGrid =(nElements + threadsPerBlock - 1) / threadsPerBlock;
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, nElements);
    //前2句，表示Grid，block都是1维时，设置网格内的块数，每块内的线程数。
    //最后一句，启动kernel（运行在gpu端的函数）函数。
    //注意前2句可以改成。dim3 threadsPerBlock(256);这种形式。

4. err = cudaGetLastError();
//启动kernel函数时，并没有返回值，通过这个调用这个函数，查看kernel函数是否启动成功。

5. err = cudaFree(d_A);
//释放使用cudaMalloc申请的空间。

6. err = cudaMemset(d_a, 0, size)
//类似于memset函数。将d_A的size个字节置0.

当gpu有好几块时，获取设备，选择设备的几个函数：

1. int num_gpus = 0;
  cudaGetDeviceCount(&num_gpus);//若num_gpus为0，表示没有检测到gpu。
2.  cudaDeviceProp dprop;
    cudaGetDeviceProperties(&dprop, id);//i代表第几块gpu，从0开始标号。
    printf("   %d: %s\n", i, dprop.name);or
    printf("Device %d: \"%s\" with Compute %d.%d capability\n",devID, props.name, props.major, props.minor);

3. err = cudaSetDevice(id)；//选择第id个设备。
    err = cudaGetDevice(&gpu_id)；//获取选择的设备号，通过id返回。

4. devID = findCudaDevice(argc, (const char **)argv);
    //devID为int类型。系统帮你选择最合适的gpu设备。此时DevID里存的是被选择的设备号吗。自己试试看看。
    err = cudaGetDevice(&devID);
    //选完之后，获取设备号。
5. error = cudaDeviceSynchronize();//TODO???
6.

关于计算地址：个人感觉这里是最迷糊人的。需要好好想清楚。
思路：先求块首地址，再加块内地址。

grid是一维，block是一维，如何计算线程号：

注意：这个好理解。blockDim相当于一个块内的线程总数。blockIdx.x相当于当前线程是第几块。threadIdx.x相当于当前线程时第几个线程（块内）。（有空时画个图便于理解）
Grid是2维，block是3维：

dim3 dimGrid(2, 2);
dim3 dimBlock(2, 2, 2);

全局块号：blockIdx.y*gridDim.x+blockIdx.x
全局块内线程号：threadIdx.z*blockDim.x*blockDim.y+threadIdx.y*blockDim.x+threadIdx.x
注意：2维的可以以列优先存储的矩阵方式来理解。3维的时候，可以理解成多层的矩阵。

关于在设备上打印(输出，主要用于debug)：

compute capability less than 2.0, the function cuPrintf is called.
others call printf();