CUDA 学习笔记 —— （一）CUDA介绍及编程流程

引出GPU并行编程

GPU计算能力这么强，被广泛使用！比如挖矿（比特币）、图形图像处理、数值模拟、机器学习算法训练等等，那我们怎么发挥GPU超强的计算能力呢？—编程！

怎么进行GPU编程呢？现在GPU形形色色，比如Nvidia、AMD、Intel都推出了自己的GPU，其中最为流行的就是Nvidia的GPU，其还推出了CUDA并行编程库。然而每个GPU生产公司都推出自己的编程库显然让学习成本上升很多，因此苹果公司就推出了标准OpenCL，说各个生产商都支持我的标准，只要有一套OpenCL的编程库就能对各类型的GPU芯片适用。当然了，OpenCL做到通用不是没有代价的，会带来一定程度的性能损失，在Nvidia的GPU上，CUDA性能明显比OpenCL高出一大截。目前CUDA和OpenCL是最主流的两个GPU编程库。

从编程语言角度看，CUDA和OpenCL都是原生支持C/C++的，其它语言想要访问还有些麻烦，比如Java，需要通过JNI来访问CUDA或者OpenCL。基于JNI，现今有各种Java版本的GPU编程库，比如JCUDA等。另一种思路就是语言还是由java来编写，通过一种工具将java转换成C。

名词概念

• ROCM：是AMD使用的平台，直接对标的是CUDA应用在NVdia上，使用的是HIP编程
• HIP：编程可以运行在ROCM或者CUDA平台上，在CUDA上会先转成CUDA编译运行，而在ROCM上直接运行。主要针对A卡
• CUDA：Nvidia 推出的CUDA平台，主要针对N卡
• OPENCL：苹果推出的标准编程库，适配各类GPU，是类似于CUDA的一个模型，但是CUDA只NVdia使用，而opencL支持所有开发商

CUDA和OpenCL的区别

一个示例

编程流程：

1. 定义kernel函数
2. 分配和初始化host上的内存数据
3. 分配和初始化device内存上的数据
4. GPU激活kernel，也就是在GPU上执行函数
5. 将kernel执行输出的结果拷贝回host
6. 清理

1. 定义kernel函数

global 定义kernel函数，kernel函数会运行在GPU上
比如：下面例子将a+b得到c

// Kernel definition
// Run on GPU
// Adding 2 numbers and store the result in c
__global__ void add(int *a, int *b, int *c)
{
	*c = *a + *b;
}

2. 分配和初始化host上的内存数据

int main(void) {
	// Allocate & initialize host data - run on the host
	int a, b, c; // host copies of a, b, c
	a = 2;
	b = 7;
}

3. 在device上分配内存，并将host数据拷贝到device

// device copies of a, b, c
int *d_a, *d_b, *d_c; 
// Allocate space for device copies of a, b, c
cudaMalloc((void **)&d_a, size);
cudaMalloc((void **)&d_b, size);
cudaMalloc((void **)&d_c, size);
// Copy a & b from the host to the device
cudaMemcpy(d_a, &a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, &b, size, cudaMemcpyHostToDevice);

4. 激活kernel函数

// Launch add() kernel on GPU with parameters (d_a, d_b, d_c)
add<<<1,1>>>(d_a, d_b, d_c);

为了提供数据并行性，每个线程是单独执行并且同时执行的。
每个block有多个线程。
这些线程的block可以被GPU上的流处理器（SM）调度。
在这个例子中我们只使用了一个block中的一个线程。

主机端调用和函数采用如下形式
kernel<<<Dgrid, Dblock>>>(param list);
Dgrid: int型或者dim3类型(x,y,z).用于定义一个grid中的block是如何组织的，可以是二维三维
Dblock: int型或者dim3类型(x,y,z). 用于定义一个block中的thread是如何组织的，也可以是二维或者三维

5. 将kernel执行输出的结果拷贝回host和清理

// Copy result back to the host
cudaMemcpy(&c, d_c, size, cudaMemcpyDeviceToHost);
// Cleanup
cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);

6. 整合代码

// Kernel definition
// Run on GPU
__global__ void add(int *a, int *b, int *c) {
    *c = *a + *b;

}
int main(void) {
    // Allocate & initialize host data - run on the host
    int a, b, c; // host copies of a, b, c
    a = 2;
    b = 7;
    int *d_a, *d_b, *d_c; // device copies of a, b, c
    // Allocate space for device copies of a, b, c
    int size = sizeof(int);
    cudaMalloc((void **)&d_a, size);
    cudaMalloc((void **)&d_b, size);
    cudaMalloc((void **)&d_c, size);
    // Copy a & b from the host to the device
    cudaMemcpy(d_a, &a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, &b, size, cudaMemcpyHostToDevice);
    // Launch add() kernel on GPU
    add<<<1,1>>>(d_a, d_b, d_c);
    // Copy result back to the host
    cudaMemcpy(&c, d_c, size, cudaMemcpyDeviceToHost);
    printf("test passed!");
    // Cleanup
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);

    return 0;
}