CUDA简单的基础概念
一、简单的实现向量加法
我们打开cuda的官网（点击这个连接地址），有个简单的流程代码：
// Device code
__global__ void VecAdd(float* A, float* B, float* C, int N)
{
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < N)
C[i] = A[i] + B[i];
}
// Host code
int main()
{
int N = ...;
size_t size = N * sizeof(float);
// Allocate input vectors h_A and h_B in host memory
float* h_A = (float*)malloc(size);
float* h_B = (float*)malloc(size);
float* h_C = (float*)malloc(size);
// Initialize input vectors
...
// Allocate vectors in device memory
float* d_A;
cudaMalloc(&d_A, size);
float* d_B;
cudaMalloc(&d_B, size);
float* d_C;
cudaMalloc(&d_C, size);
// Copy vectors from host memory to device memory
cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
// Invoke kernel
int threadsPerBlock = 256;
int blocksPerGrid =
(N + threadsPerBlock - 1) / threadsPerBlock;
VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
// Copy result from device memory to host memory
// h_C contains the result in host memory
cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
// Free device memory
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
// Free host memory
...
}我们翻译一下各个步骤你就可以直接看出来了使用GPU进行向量加法的步骤了：
// vector_add.cu
#include <stdio.h>
// CUDA核函数，用于在GPU上执行向量加法
__global__ void vectorAdd(int *a, int *b, int *c, int N) {
// 获取当前线程的索引
int i = blockIdx.x * blockDim.x + threadIdx.x;
// 确保线程索引在向量大小范围内
if (i < size) {
// 计算向量元素相加结果
c[i] = a[i] + b[i];
}
}
int main() {
int size = 1000;
int a[size], b[size], c[size];
// 在设备上分配内存空间
int *dev_a, *dev_b, *dev_c;
cudaMalloc((void**)&dev_a, size * sizeof(int));
cudaMalloc((void**)&dev_b, size * sizeof(int));
cudaMalloc((void**)&dev_c, size * sizeof(int));
// 初始化输入向量 a 和 b
for (int i = 0; i < size; ++i) {
a[i] = i;
b[i] = i * 2;
}
// 将输入向量 a 和 b 复制到设备内存
cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
// 定义每个块中的线程数和块数
int threadsPerBlock = 256;
int blocksPerGrid = (size + threadsPerBlock - 1) / threadsPerBlock;
// 调用 CUDA 核函数执行向量加法
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(dev_a, dev_b, dev_c, size);
// 将计算结果从设备复制到主机内存
cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
// 打印结果向量 c 的前10个元素
for (int i = 0; i < 10; ++i) {
printf("%d ", c[i]);
}
printf("\n");
// 释放设备上的内存空间
cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);
return 0;
}从上面的注释我们就可以看出CUDA基础开发流程是什么样子了。
二、CUDA基础开发流程：
在CUDA编程中，设备端的大部分程序代码可以独立于主机端的程序代码执行，当一个Kernel函数被调用后，控制权会返回到主机端的CPU控制。一般来说，CPU负责串行的工作而由GPU进行并行的操作。
CUDA的开发流程虽然因为不同的接口方式有所不同，但一般分以下几步：
1、分配CPU和GPU内存
2、复制数据从CPU端到GPU端
3、调用CUDA kernel来执行运算
4、运算结束将结果数据从GPU拷贝至CPU
5、释放CPU和GPU内存
三、基础概念
通过代码展示一个简单的CUDA实现向量加法的流程我们可以看到里面涉及到了很多的概念：
1、CPU和GPU
CPU我们可以把它看做一个指挥者，主机端，host，而完成大量计算的GPU是我们的计算设备，device。
- 左图：四核CPU=4个ALU（算数逻辑单元）+1个控制单元（control）+1个缓存（cache），内存（DRAM）一般不在片上，CPU通过总线访问内存。
- 右图：GPU，N个ALU，红色框SM内所有ALU公用一个Control单元和Cache
2、SM
一个SM相当于一个完整的多核CPU，但是ALU多了（计算能力提升），control少了（控制能力减弱），所以对于控制（逻辑）复杂的程序，GPU的SM没法和CPU比较，但是对于逻辑简单，数据量大的任务，GPU更高效。并且，一个GPU有非常多SM且越来越多。
在 GPU 中，SM 通常是指 “Streaming Multiprocessor”，即流式多处理器。它是 NVIDIA GPU 架构中的核心组件，相当于 GPU 的 “心脏”，类似于 CPU 的核心，负责执行图形和计算操作。
每个 SM 包含多个流处理器（CUDA Core）或线程，还配备了共享内存、控制单元和其他加速计算单元，如张量核心或光线追踪单元等。SM 是 GPU 进行资源分配和调度的基本单位，GPU 会把线程块分配到 SM 上，由 SM 负责调度块内的线程执行计算任务。SM 的数量和性能是衡量 GPU 整体性能的重要指标之一，其数量越多，GPU 的并行计算能力通常越强。
3、CUDA基础框架说明
CUDA编程中，是通过CPU和GPU共同来完成的，CPU居于主导的控制地位而GPU用于协助进行大规模的数据计算和处理。CPU与GPU通过PCIE总线连接到一起，GPU之间通信通过PCIe或者是NVLink来实现。CPU被称为主机端（host）而GPU端被称为设备端（device）
4、 CUDA编程结构
异构环境中的多个CPU和多个GPU通过PCIe总线相互通信，也通过PCIe总线相互分隔，需要区分两种设备的内存：
• 主机：CPU及其内存
• 设备：GPU及其内存
这两个内存从硬件到软件都是隔离的，需要用内存来回拷贝的方法编写程序
核函数被调用后控制立刻归还给主机线程，即第一个并行代码执行时，很有可能第二段host代码已经开始同步执行了
5、CUDA编程模型概述
编程模型可以理解为：用于控制异构计算设备工作模式的语法，内存结构，线程结构等，GPU中大致可以分为：
- 核函数
- 内存管理
- 线程管理
- 流
（1）、核函数
CUDA 核函数（Kernel Function）是 CUDA 编程模型的核心，它定义了在 GPU 上并行执行的计算任务。
简单来说，你可以将一个 CUDA 程序理解为由两部分组成：
- 主机（Host）代码：运行在 CPU 上，负责处理串行逻辑、数据准备和启动核函数。
- 设备（Device）代码：即核函数，运行在 GPU 上，负责处理大规模并行计算任务。
CUDA 核函数在 C/C++ 代码中定义，通过 __global__ 关键字声明。它与普通函数有明显区别，普通函数在 CPU 上顺序执行，而核函数在 GPU 的多个线程上并行执行。例如：
__global__ void vectorAdd(float* a, float* b, float* c, int n) {
// 计算当前线程的全局索引
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}核函数调用时会启动大量线程并行执行。每个线程执行相同的核函数代码，但处理不同的数据。这些线程以层次化的结构组织，包括线程块（block）和网格（grid）。上述 addKernel 核函数，每个线程根据自身索引 index 从输入数组 a 和 b 中取出对应元素相加，结果存入数组 c
核函数不能像普通函数那样直接调用，必须通过一个特殊的语法 <<< >>> 来配置其并行执行的规模。
// 语法：核函数名<<<网格配置, 线程块配置>>>(参数列表); vectorAdd<<<numBlocks, threadsPerBlock>>>(d_a, d_b, d_c, N);
这里的 <<< >>> 被称为执行配置，它定义了线程的层次结构
线程（Thread）：GPU 执行的最小单位。每个线程独立执行核函数中的代码。
线程块（Block）：一组线程的集合。块内的线程可以通过共享内存（Shared Memory）快速通信，并使用 __syncthreads() 进行同步。
网格（Grid）：包含所有线程块的集合。 例如，<<<10, 256>>> 表示启动一个包含 10 个线程块的网格，每个块内有 256 个线程，总共将有 2560 个线程并行执行该核函数。
（2）、内存管理
CUDA的API可以分配管理Device上的内存，当然也可以用CDUA管理Host上的内存，主机上的传统标准库也能完成主机内存管理
内存分层结构如下图所示：
（3）、线程管理
a、分层组织结构
一个核函数只能有一个grid，一个grid可以有很多个块，每个块可以有很多的线程：
同一个线程块block中的线程可以完成下述协作：
- 同步
- 共享内存
==不同线程块block内线程不能相互影响！他们是物理隔离的！==
b、线程编号
每个线程执行同样的一段串行代码，如何让同一段代码对应不同的数据？
首先需要将线程彼此区分开，使得这些线程也能区分自己的数据：
- blockIdx（block在grid内的位置索引，block index）
- threadIdx（thread在block内的位置索引，thread index）
因为一个核函数只有一个grid，所以不需要gridIdx
这两个内置结构体基于uint3定义，包含3个无符号整数的结构，通过3个字段指定：
①blockIdx
- blockIdx.x // block在grid中的x坐标
- blockIdx.y // block在grid中的y坐标
- blockIdx.z // block在grid中的z坐标
gridDim表示grid中block索引的范围，grid一般是二维的即grid通常被分为2D blcok
②threadIdx
- threadIdx.x // thread在block中的x坐标
- threadIdx.y // thread在block中的y坐标
- threadIdx.z // thread在block中的z坐标
blockDim表示block中thread索引的范围，block一般是三维的即block通常被分为3D thread
blockDim.x、blockDim.y、blockDim.z
注意：dim3是手工定义的，主机端可见。uint3是设备端在执行的时候可见的，不可以在核函数运行时修改，初始化完成后uint3值就不变了。
6、CUDA重要的关键字
在CUDA编程中主要的关键字有：
1、__global__
此关键字用来定义核函数，由<<< >>>配置在执行核心的执行线程数，每个线程都有唯一一个线程ID并可在内核中进行访问。它的特点包括：
主机端或计算能力在大于等于3.2的设备端调用，在设备端执行；
其返回值必须为void类型且不能是类成员函数；
必须指定其执行配置即上面提到的<<<>>>;
函数的调用是异步的即其在设备执行完成前返回；
2、__device__
它被定义为只能由设备端调用并且在设备端执行的函数，需要注意的是，它不能与关键字__global__一起使用
3、__host__
这个关键字用来定义在主机上执行的函数，其特点为：
函数只能在主机端调用和执行；
它不可以和关键字__global__一起使用，但可以与__device__一起使用
本文原创作者：冯一川（csdn：ifeng12358），未经作者授权同意，请勿转载。
扫一扫
2183
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
