CUDA 是一个并行计算框架. 用于计算加速. 是 nvidia 家的产品. 广泛地应用于现在的深度学习加速.
一句话描述就是: cuda 帮助我们把运算从 CPU 放到 gpu 上做, gpu 多线程同时处理运算, 达到加速效果.
从一个简单例子说起:
- #include <iostream>
- #include <math.h>
- // function to add the elements of two arrays
- void add(int n, float *x, float *y)
- {
- for (int i = 0; i <n; i++)
- y[i] = x[i] + y[i];
- }
- int main(void)
- {
- int N = 1<<20; // 1M elements
- float *x = new float[N];
- float *y = new float[N];
- // initialize x and y arrays on the host
- for (int i = 0; i < N; i++) {
- x[i] = 1.0f;
- y[i] = 2.0f;
- }
- // Run kernel on 1M elements on the CPU
- add(N, x, y);
- // Check for errors (all values should be 3.0f)
- float maxError = 0.0f;
- for (int i = 0; i < N; i++)
- maxError = fmax(maxError, fabs(y[i]-3.0f));
- std::cout << "Max error:" << maxError << std::endl;
- // Free memory
- delete [] x;
- delete [] y;
- return 0;
- }
这段代码很简单, 对两个数组对应位置元素相加. 数组很大, 有 100 万个元素.
代码运行时间在 0.075s.
改写代码使之运行于 gpu
gpu 上能够运算的函数, 在 cuda 中我们称之为 kernel.由 nvcc 将其编译为可以在 GPU 上运行的格式.
- #include <iostream>
- #include <math.h>
- // Kernel function to add the elements of two arrays
- __global__
- void add(int n, float *x, float *y)
- {
- for (int i = 0; i <n; i++)
- y[i] = x[i] + y[i];
- }
- int main(void)
- {
- int N = 1<<20;
- float *x, *y;
- // Allocate Unified Memory - accessible from CPU or GPU
- cudaMallocManaged(&x, N*sizeof(float));
- cudaMallocManaged(&y, N*sizeof(float));
- // initialize x and y arrays on the host
- for (int i = 0; i < N; i++) {
- x[i] = 1.0f;
- y[i] = 2.0f;
- }
- // Run kernel on 1M elements on the GPU
- add<<<1, 1>>>(N, x, y);
- // Wait for GPU to finish before accessing on host
- cudaDeviceSynchronize();
- // Check for errors (all values should be 3.0f)
- float maxError = 0.0f;
- for (int i = 0; i <N; i++)
- maxError = fmax(maxError, fabs(y[i]-3.0f));
- std::cout << "Max error:" << maxError << std::endl;
- // Free memory
- cudaFree(x);
- cudaFree(y);
- return 0;
- }
nvcc 编译的文件的后缀为. cu
cuda 中定义 kernel 在函数前加上__global 声明就可以了.
在显存上分配内存使用 cudaMallocManaged
调用一个函数使用 <<<>>> 符号. 比如对 add 的函数的调用使用`add<<<1, 1>>>(N, x, y);`, 关于其中参数的意义, 后文再做解释.
需要 cudaDeviceSynchronize()让 CPU 等待 gpu 上的计算做完再执行 CPU 上的操作
可以用 nvprof 做更详细的性能分析.
注意用 sudo 否则可能报错.
sudo /usr/local/cuda/bin/nvprof ./add_cuda
gpu 上 add 用了 194ms.
这里, 我们注意到, 跑在 gpu 反而比 CPU 更慢了. 因为我们这段代码里`add<<<1, 1>>>(N, x, y);`并没有发挥 gpu 并行运算的优势, 反而因为多了一些 CPU 与 gpu 的交互使得程序变慢了.
用 GPU threads 加速运算
重点来了
CUDA GPUS 有多组 Streaming Multiprocessor(SM). 每个 SM 可以运行多个 thread block. 每一个 thread block 有多个 thread.
如下图所示:
注意几个关键变量:
blockDim.x 表明了一个 thread block 内含有多少个 thread
threadIdx.x 表明了当前 thread 在该 thread blcok 内的 index
blockIdx.x 表明了当前是第几个 thread block
我们要做的就是把计算分配到所有的 thread 上去. 这些 thread 上并行地做运算, 从而达到加速的目的.
前面我们说到在 cuda 内调用一个函数 (称之为 kernel) 的用法为<<<p1,p2>>>, 比如`add<<<1, 1>>>(N, x, y);` 第一个参数的含义即为 thread block 的数量, 第二个参数的含义为 block 内参与运算的 thread 数量.
现在来改写一下代码:
- #include <iostream>
- #include <math.h>
- #include <stdio.h>
- // Kernel function to add the elements of two arrays
- __global__
- void add(int n, float *x, float *y)
- {
- int index = threadIdx.x;
- int stride = blockDim.x;
- printf("index=%d,stride=%d\n",index,stride);
- for (int i = index; i <n; i+=stride)
- {
- y[i] = x[i] + y[i];
- if(index == 0)
- {
- printf("i=%d,blockIdx.x=%d,thread.x=%d\n",i,blockIdx.x,threadIdx.x);
- }
- }
- }
- int main(void)
- {
- int N = 1<<20;
- float *x, *y;
- // Allocate Unified Memory - accessible from CPU or GPU
- cudaMallocManaged(&x, N*sizeof(float));
- cudaMallocManaged(&y, N*sizeof(float));
- // initialize x and y arrays on the host
- for (int i = 0; i < N; i++) {
- x[i] = 1.0f;
- y[i] = 2.0f;
- }
- // Run kernel on 1M elements on the GPU
- add<<<1, 256>>>(N, x, y);
- // Wait for GPU to finish before accessing on host
- cudaDeviceSynchronize();
- // Check for errors (all values should be 3.0f)
- float maxError = 0.0f;
- for (int i = 0; i <N; i++)
- maxError = fmax(maxError, fabs(y[i]-3.0f));
- std::cout << "Max error:" << maxError << std::endl;
- // Free memory
- cudaFree(x);
- cudaFree(y);
- return 0;
- }
注意 add 的写法, 我们把 0,256,512... 放到 thread1 计算, 把 1,257,... 放到 thread2 计算, 依次类推. 调用的时候,
add<<<1, 256>>>(N, x, y);
表明我们只把计算分配到了 thread block1 内的 256 个 thread 去做.
编译这个程序(注意把代码里的 printf 注释掉, 因为要统计程序运行时间):nvcc add_block.cu -o add_cuda_blcok -I/usr/local/cuda-9.0/include/ -L/usr/local/cuda-9.0/lib64
可以看到 add 的 gpu 时间仅仅用了 2.87ms
程序的整体运行时间为 0.13s, 主要是 cudaMallocManaged,cudaDeviceSynchronize 之类的操作耗费了比较多的时间.
再一次改写代码
这一次我们用更多的 thread block.
- int blockSize = 256;
- int numBlocks = (N + blockSize - 1) / blockSize;
- add<<<numBlocks, blockSize>>>(N, x, y);
- // Kernel function to add the elements of two arrays
- __global__
- void add(int n, float *x, float *y)
- {
- int index = blockIdx.x * blockDim.x + threadIdx.x;
- int stride = blockDim.x * gridDim.x;
- for (int i = index; i < n; i+=stride)
- {
- y[i] = x[i] + y[i];
- //printf("i=%d,blockIdx.x=%d\n",i,blockIdx.x);
- }
- }
编译: nvcc add_grid.cu -o add_cuda_grid -I/usr/local/cuda-9.0/include/ -L/usr/local/cuda-9.0/lib64
统计性能:
可以看出来, gpu 上 add 所用的时间进一步缩小到 1.8ms
参考:
来源: https://www.cnblogs.com/sdu20112013/p/12629478.html