总说
主要是 https://devblogs.nvidia.com/even-easier-introduction-cuda/ 的简要信息提取
准备工作
我们写的cuda程序需要nvcc来进行编译, 而其实nvcc就是在安装好的cuda目录的bin文件夹下:
我们首先将相应目录添加进.bashrc
里面.
# CUDA PATH
export PATH="/usr/local/cuda-9.2/bin:$PATH"
export LD_LIBRARY_PATH="/usr/local/cuda-9.2/lib64:$LD_LIBRARY_PATH"
source .bashrc
进行更新.
接下来进入正题.
C++程序与CUDA程序的对比
新建add.cpp
#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;
}
对应的CUDA程序是add.cu
#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;
}
几个知识点:
- 我们在
add
这个函数上面加了__global__
分类符, 表示CUDA C++编译器这个函数可以在GPU上跑, 而且是由CPU进行调用的. cydaMallocManaged()
函数可以申请Unified Memory
, 而在这个统一内存可以被GPU和CPU都可以使用.cudaFree()
用来释放内存cudaDeviceSynchronize()
是用来同步一下, 即等待gpu计算完成. 再执行后面的代码add<<<1, 1>>>(N, x, y)
这个是什么?
<<<1, 1>>>
这个是什么?
简单来说就是, 要调用CUDA核, 就要用<<< >>>
来进行. 那么这两个1
是什么东西呢? 其实表示的是用多少个并行线程来进行处理.
第一个1表示,用一个线程block, 第二个1表示每个线程block使用一个线程.
cuda程序的编译以及性能测试
nvcc add.cu -o add_cuda
./add_cuda
Max error: 0.000000
现在这里只有用一个线程进行实验, 我们可以看到, 因为for
循环中, 是对整个数组都逐个进行相加, 如果我们直接用多个线程来做, 就相当于每个线程都会对所有的元素进行操作, 就会造成 race condition
的情况.
nvprof ./add_cuda
==3355== NVPROF is profiling process 3355, command: ./add_cuda
Max error: 0
==3355== Profiling application: ./add_cuda
==3355== Profiling result:
Time(%) Time Calls Avg Min Max Name
100.00% 463.25ms 1 463.25ms 463.25ms 463.25ms add(int, float*, float*)
现在我们要进行用多个线程进行操作, 需要改写一下for循环即可, 从而让每个线程处理不同的数据
// 假设一个block我们有256个线程
// 那么我们需要numBlocks个blocks
int blockSize = 256;
int numBlocks = (N + blockSize - 1) / blockSize;
add<<<numBlocks, blockSize>>>(N, x, y);
__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];
}
没错, 这是通用的写法,
我们来看看CUDA并行分层
enmm, 一个block里面有多个threads, 那么多个block组成一个grid. 没啥毛病.
gridDim.x
: 一个grid的dim是多少, 即包含多少个blockgridIdx.x
: 第几个gridblockDim.x
: 一个block包含多少个threadsblockIdx.x
: 这是第几个blockthreadIdx.x
: 第几个thread
所以再看看刚才的,
// 下面三句可以认为是定式
int index = blockIdx.x * blockDim.x + threadIdx.x;
int stride = blockDim.x * gridDim.x;
for (int i = index; i < n; i += stride)
完全没毛病, 值得注意的是, 怎么理解这个?
因为我也是初学, 所以它可能是这样工作的.
首先, 因为外面调用这个kernel是 <<<nBlocks, nTheads>>>
来的. 那么表示一下子会用nBlocks * nThreads
这么多个线程进行工作. 那么, 对于单个线程进入这个函数时, 函数内部通过 index
来获得这个线程所处理的元素是x[i]
, 其中i=index
. 显然啊, 这个index
是看这个线程所在的block
属于我们申请的所有的block中的第几个, 以及这个线程是block中的第几个线程. 因此index = blockIdx.x * blockDim.x + threadIdx.x
.
申请的所有的线程同时进行相同的工作, 而for循环中的index恰好可以把每个线程处理哪个元素进行很好的对应分配, 因此经过一次运算时间, 有nBlocks * nThreads
个元素已经处理完了. 自然我们要转向下一组数据, 即此时有int stride = blockDim.x * gridDim.x
以及i = i+stride
.
从代码上看, 这个blockIdx.x
很有可能是从0
(申请的blocks的第一个)开始的, 不过可能也不一定, 这个以后更多理解之后, 再看看是否正确!
Time(%) Time Calls Avg Min Max Name
100.00% 94.015us 1 94.015us 94.015us 94.015us add(int, float*, float*)
到现在为止, 我们就可以告一段落.
总结
简单来说, 知道了__global__
分类符是用来修饰函数, 表示该函数(称为kernel)可以被GPU调用. 然后还知道了 CUDA 的for循环的编程3句定式, 还知道了CUDA申请内存以及释放的相关函数等.