文章目录
参考文献
CUDA编程结构
异构环境中，包含多个CPU以及GPU，他们之间通过PCIE总线互相通信，也是通过PCIE总线分隔开
- 主机：cpu及其内存
- 设备：GPU及其内存
注意这两个内存从硬件到软件都是隔离的（CUDA6.0 以后支持统一寻址）
CUDA应用代码执行顺序
- 这里Host code执行结束之后，调用核函数开始执行parallel code。
- GPU在执行parallel code的时候，CPU继续向下执行Host code
内存管理
CUDA提供API来管理device上的内存，也有API可以管理Host上的内存
| 标准C函数 | CUDA C函数 | 说明 |
|---|---|---|
| malloc | cudaMalloc | 内存分配 |
| memcpy | cudaMemcpy | 内存复制 |
| memset | cudaMemset | 内存设置 |
| free | cudaFree | 释放内存 |
在cuda中，内存是分层次的，这里不进行详细讲，只是大概了解一下
一定注意区分Host 内存和 device 内存
补充：四种函数的用法
cudaMalloc、cudaMemcpy、cudaMemset 和 cudaFree 是 CUDA 编程中用于设备内存管理的核心函数。它们分别用于分配设备内存、在主机和设备之间传输数据、初始化设备内存以及释放设备内存。
-
cudaMalloc:
- 功能：在设备上分配指定大小的内存。
- 函数原型：
cudaError_t cudaMalloc(void** devPtr, size_t size) - 参数：
devPtr: 指向设备内存指针的指针，分配的内存地址将存储在这里。size: 要分配的内存大小，以字节为单位。
- 返回值：返回
cudaSuccess表示成功，否则返回错误代码。 - 示例：
float* d_array; cudaMalloc((void**)&d_array, 100 * sizeof(float));
-
cudaMemcpy:
- 功能：在主机和设备之间或设备与设备之间复制数据。
- 函数原型：
cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind) - 参数：
dst: 目标内存地址。src: 源内存地址。count: 要复制的字节数。kind: 复制方向，可以是cudaMemcpyHostToHost、cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost或cudaMemcpyDeviceToDevice。
- 返回值：返回
cudaSuccess表示成功，否则返回错误代码。 - 示例：
float h_array[100]; float* d_array; cudaMalloc((void**)&d_array, 100 * sizeof(float)); cudaMemcpy(d_array, h_array, 100 * sizeof(float), cudaMemcpyHostToDevice);
-
cudaMemset:
- 功能：将设备内存的指定区域设置为特定的值。
- 函数原型：
cudaError_t cudaMemset(void* devPtr, int value, size_t count) - 参数：
devPtr: 设备内存的起始地址。value: 要设置的值（以int形式传递，但实际按字节设置）。count: 要设置的字节数。
- 返回值：返回
cudaSuccess表示成功，否则返回错误代码。 - 示例：
int* d_array; cudaMalloc((void**)&d_array, 100 * sizeof(int)); cudaMemset(d_array, 0, 100 * sizeof(int));
-
cudaFree:
- 功能：释放之前通过
cudaMalloc分配的设备内存。 - 函数原型：
cudaError_t cudaFree(void* devPtr) - 参数：
devPtr: 要释放的设备内存指针。
- 返回值：返回
cudaSuccess表示成功，否则返回错误代码。 - 示例：
float* d_array; cudaMalloc((void**)&d_array, 100 * sizeof(float)); cudaFree(d_array);
- 功能：释放之前通过
-
案例
#include <iostream>
#include <cuda_runtime.h>
// CUDA 错误检查宏
#define CHECK_CUDA(call) \
do { \
cudaError_t err = call; \
if (err != cudaSuccess) { \
std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__ << " - " << cudaGetErrorString(err) << std::endl; \
exit(EXIT_FAILURE); \
} \
} while (0)
int main() {
const int N = 10; // 数组大小
size_t size = N * sizeof(int); // 计算内存大小
// 主机内存分配
int* h_array = new int[N];
// 初始化主机数组
for (int i = 0; i < N; i++) {
h_array[i] = i;
}
// 设备内存分配
int* d_array;
CHECK_CUDA(cudaMalloc((void**)&d_array, size));
// 将主机数据复制到设备
CHECK_CUDA(cudaMemcpy(d_array, h_array, size, cudaMemcpyHostToDevice));
// 在设备上初始化内存
int value = 0;
CHECK_CUDA(cudaMemset(d_array, value, size));
// 将设备数据复制回主机
CHECK_CUDA(cudaMemcpy(h_array, d_array, size, cudaMemcpyDeviceToHost));
// 打印结果
std::cout << "Host array after cudaMemset: ";
for (int i = 0; i < N; i++) {
std::cout << h_array[i] << " ";
}
std::cout << std::endl;
// 释放设备内存
CHECK_CUDA(cudaFree(d_array));
// 释放主机内存
delete[] h_array;
return 0;
}
线程管理
当内核函数开始执行，最重要的事情就是如何组织GPU的线程
GPU线程按层次组织，通常分为三个级别：
- 线程（Thread）：最基本的执行单元。
- 线程块（Block）：包含多个线程，线程块内的线程可以同步和共享内存。
- 网格（Grid）：包含多个线程块，网格中的所有线程块执行相同的核函数。
一个核函数只能有一个grid，一个grid 会分成很多个块，每个块可以分成很多个线程
- 分成块、线程，可以是一维、二维、三维的，一般而言，一个grid会被分成二维的block，一个block会被分成三维的线程
- 一个block中的线程，可以完成
同步和共享内存 - 不同block中的线程不能互相影响，他们是物理隔离的
线程编号
在GPU编程中，线程的编号和层次结构是通过一组内置变量来确定的。这些变量帮助开发者明确每个线程的唯一标识，从而在并行计算中分配任务。
threadIdx：线程在线程块内的编号
- 类型：
dim3（三维向量，包含x、y、z三个分量）。 - 含义：表示当前线程在其所属线程块中的局部编号。
- 范围：
threadIdx.x：范围是[0, blockDim.x - 1]。threadIdx.y：范围是[0, blockDim.y - 1]。threadIdx.z：范围是[0, blockDim.z - 1]。
- 用途：用于标识线程在线程块中的位置，通常用于处理多维数据（如图像、矩阵等）。
blockIdx：线程块在网格中的编号
- 类型：
dim3（三维向量，包含x、y、z三个分量）。 - 含义：表示当前线程块在其所属网格中的编号。
- 范围：
blockIdx.x：范围是[0, gridDim.x - 1]。blockIdx.y：范围是[0, gridDim.y - 1]。blockIdx.z：范围是[0, gridDim.z - 1]。
- 用途：用于标识线程块在网格中的位置，通常用于处理更大的数据集或任务分解。
blockDim：线程块的维度
- 类型：
dim3（三维向量，包含x、y、z三个分量）。 - 含义：表示当前线程块的维度，即线程块中包含的线程数量。
- 范围：
blockDim.x：线程块在x方向的线程数。blockDim.y：线程块在y方向的线程数。blockDim.z：线程块在z方向的线程数。
- 用途：用于计算线程的全局编号，以及确定线程块的大小。
gridDim：网格的维度
- 类型：
dim3（三维向量，包含x、y、z三个分量）。 - 含义：表示当前网格的维度，即网格中包含的线程块数量。
- 范围：
gridDim.x：网格在x方向的线程块数。gridDim.y：网格在y方向的线程块数。gridDim.z：网格在z方向的线程块数。
- 用途：用于计算线程块的全局编号，以及确定网格的大小。
- 线程的全局编号计算
线程的全局编号可以通过以下公式计算：
- 一维网格和线程块：
int globalThreadId = blockIdx.x * blockDim.x + threadIdx.x; - 二维网格和线程块：
int globalThreadIdX = blockIdx.x * blockDim.x + threadIdx.x; int globalThreadIdY = blockIdx.y * blockDim.y + threadIdx.y; - 三维网格和线程块：
int globalThreadIdX = blockIdx.x * blockDim.x + threadIdx.x; int globalThreadIdY = blockIdx.y * blockDim.y + threadIdx.y; int globalThreadIdZ = blockIdx.z * blockDim.z + threadIdx.z;
在GPU编程中，网格（Grid）和线程块（Block）可以定义为一维、二维或三维结构，具体取决于问题的需求。以下是 一维、二维 和 三维 网格和线程块的定义示例，以及如何在核函数中使用它们。
一维网格和线程块
定义：
// 定义一维网格和线程块
// 可以用dim3
int numBlocks = 4; // 网格中有 4 个线程块
int threadsPerBlock = 8; // 每个线程块有 8 个线程
// 启动核函数
myKernel<<<numBlocks, threadsPerBlock>>>();
核函数中的线程编号：
__global__ void myKernel() {
int threadId = blockIdx.x * blockDim.x + threadIdx.x;
printf("Block %d, Thread %d, Global Thread ID %d\n", blockIdx.x, threadIdx.x, threadId);
}
二维网格和线程块
定义：
// 定义二维网格和线程块
dim3 gridDim(2, 3); // 网格有 2x3 个线程块
dim3 blockDim(4, 4); // 每个线程块有 4x4 个线程
// 启动核函数
myKernel<<<gridDim, blockDim>>>();
核函数中的线程编号：
__global__ void myKernel() {
int threadIdX = blockIdx.x * blockDim.x + threadIdx.x;
int threadIdY = blockIdx.y * blockDim.y + threadIdx.y;
printf("Block (%d, %d), Thread (%d, %d), Global Thread ID (%d, %d)\n",
blockIdx.x, blockIdx.y, threadIdx.x, threadIdx.y, threadIdX, threadIdY);
}
三维网格和线程块
定义：
// 定义三维网格和线程块
dim3 gridDim(2, 2, 2); // 网格有 2x2x2 个线程块
dim3 blockDim(3, 3, 3); // 每个线程块有 3x3x3 个线程
// 启动核函数
myKernel<<<gridDim, blockDim>>>();
核函数中的线程编号：
__global__ void myKernel() {
int threadIdX = blockIdx.x * blockDim.x + threadIdx.x;
int threadIdY = blockIdx.y * blockDim.y + threadIdx.y;
int threadIdZ = blockIdx.z * blockDim.z + threadIdx.z;
printf("Block (%d, %d, %d), Thread (%d, %d, %d), Global Thread ID (%d, %d, %d)\n",
blockIdx.x, blockIdx.y, blockIdx.z, threadIdx.x, threadIdx.y, threadIdx.z,
threadIdX, threadIdY, threadIdZ);
}
完整示例
以下是一个完整的二维网格和线程块的示例：
#include <stdio.h>
__global__ void myKernel() {
int threadIdX = blockIdx.x * blockDim.x + threadIdx.x;
int threadIdY = blockIdx.y * blockDim.y + threadIdx.y;
printf("Block (%d, %d), Thread (%d, %d), Global Thread ID (%d, %d)\n",
blockIdx.x, blockIdx.y, threadIdx.x, threadIdx.y, threadIdX, threadIdY);
}
int main() {
// 定义二维网格和线程块
dim3 gridDim(2, 2); // 2x2 网格
dim3 blockDim(3, 3); // 3x3 线程块
// 启动核函数
myKernel<<<gridDim, blockDim>>>();
cudaDeviceSynchronize(); // 等待 GPU 完成
return 0;
}
输出：
Block (0, 0), Thread (0, 0), Global Thread ID (0, 0)
Block (0, 0), Thread (0, 1), Global Thread ID (0, 1)
Block (0, 0), Thread (0, 2), Global Thread ID (0, 2)
Block (0, 0), Thread (1, 0), Global Thread ID (1, 0)
Block (0, 0), Thread (1, 1), Global Thread ID (1, 1)
Block (0, 0), Thread (1, 2), Global Thread ID (1, 2)
Block (0, 0), Thread (2, 0), Global Thread ID (2, 0)
Block (0, 0), Thread (2, 1), Global Thread ID (2, 1)
Block (0, 0), Thread (2, 2), Global Thread ID (2, 2)
Block (0, 1), Thread (0, 0), Global Thread ID (0, 3)
Block (0, 1), Thread (0, 1), Global Thread ID (0, 4)
...
Block (1, 1), Thread (2, 2), Global Thread ID (5, 5)
核函数 Kernel Function
在GPU编程中，核函数（Kernel Function） 是一个在GPU上并行执行的函数。它是GPU编程的核心部分，用于定义在大量线程上并行执行的任务。核函数由CPU调用，但在GPU上执行，利用GPU的并行计算能力来加速计算任务。
核函数的特点
- 并行执行：核函数在成千上万的线程上同时执行，每个线程独立运行相同的代码，但可以处理不同的数据。
- 由CPU调用：核函数由CPU通过特定的语法启动，但实际执行在GPU上。
- 线程层次结构：核函数中的线程被组织成线程块（Block）和网格（Grid），开发者可以通过内置变量（如 threadIdx、blockIdx 等）来标识每个线程的唯一编号。
- 无返回值：核函数通常没有返回值，其计算结果通过全局内存、共享内存或直接输出到设备内存中。
核函数的定义
CUDA C中的函数限定符
CUDA C扩展了C语言，引入了一些新的函数限定符，用于指定函数的执行位置和调用方式。以下是这些限定符的说明：
| 限定符 | 执行位置 | 调用方式 | 备注 |
|---|---|---|---|
__global__ | 设备端执行 | 可以从主机调用，也可以从计算能力3以上的设备调用 | 必须返回 void，用于定义核函数。 |
__device__ | 设备端执行 | 只能从设备端调用,不能被cpu调用 | 用于定义在GPU上执行的辅助函数。 |
__host__ | 主机端执行 | 只能从主机调用 | 与普通C函数一致，可以省略（默认就是 __host__）。 |
特殊情况下：同时定义
__device__和__host__
有些函数可以同时定义为 __device__ 和 __host__，这种函数可以同时被设备端和主机端的代码调用。例如：
__host__ __device__ int add(int a, int b) {
return a + b;
}
- 主机端调用：函数会被编译为CPU机器码，由CPU执行。
- 设备端调用：函数会被编译为GPU机器码，由GPU执行。
- 编译器行为：编译器会生成两份机器码，一份用于CPU，一份用于GPU。
Kernel核函数编写的限制
核函数（__global__ 函数）在编写时有以下限制：
-
只能访问设备内存
核函数只能访问GPU的设备内存（如全局内存、共享内存等），不能直接访问主机内存（CPU内存）。如果需要使用主机数据，必须先将数据拷贝到设备内存中。 -
必须有
void返回类型
核函数不能有返回值，必须返回void。计算结果通常通过修改设备内存中的值来传递。 -
不支持可变数量的参数
核函数不支持像C语言中的printf那样的可变参数列表（即...语法）。参数的数量和类型必须在编译时确定。 -
不支持静态变量
核函数中不能使用静态变量（static变量）。静态变量的生命周期和作用域与CUDA的并行执行模型不兼容。 -
显示异步行为
核函数的执行是异步的。当CPU调用核函数时，CPU会立即继续执行后续代码，而不会等待核函数执行完成。如果需要同步，必须显式调用cudaDeviceSynchronize()来等待核函数执行完毕。
核函数的调用
核函数通过 三重尖括号语法 <<<>>> 调用，用于指定网格（Grid）和线程块（Block）的维度。
__global__ void myKernel(int *data, int N) {
// 核函数代码
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
data[idx] = data[idx] * 2; // 示例：将数组中的每个元素乘以 2
}
}
int main() {
int N = 100;
int *d_data;
cudaMalloc(&d_data, N * sizeof(int)); // 在GPU上分配内存
// 定义网格和线程块
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
// 调用核函数
myKernel<<<blocksPerGrid, threadsPerBlock>>>(d_data, N);
// 将结果从GPU拷贝回CPU
int h_data[N];
cudaMemcpy(h_data, d_data, N * sizeof(int), cudaMemcpyDeviceToHost);
// 释放GPU内存
cudaFree(d_data);
return 0;
}
是否等待核函数完成
主机启动核函数后，默认不会等待device完成kernel function
如果希望等待，可以用以下两种方法
- 显式等待
cudaError_t cudaDeviceSynchronize(void);
- 隐式等待:隐式方法就是不明确说明主机要等待设备端，而是设备端不执行完，主机没办法进行，比如内存拷贝函数
cudaError_t cudaMemcpy(void* dst,const void * src,
size_t count,cudaMemcpyKind kind);
补充：__global__和__device__的区别
在CUDA编程中，__device__ 和 __global__ 是两个重要的函数限定符，它们用于定义在GPU上执行的函数，但它们的用途和行为有显著区别。以下是它们的详细对比：
1. 定义与用途
| 限定符 | 定义与用途 |
|---|---|
__global__ | 用于定义 核函数（Kernel Function），是GPU程序的入口点，由CPU调用并在GPU上并行执行。 |
__device__ | 用于定义 设备函数，是GPU内部的辅助函数，只能由其他设备函数或核函数调用。 |
2. 调用方式
| 限定符 | 调用方式 |
|---|---|
__global__ | 由CPU调用，使用 <<<>>> 语法启动。也可以从支持动态并行的GPU设备调用（计算能力3.0以上）。 |
__device__ | 只能由GPU上的其他函数（如核函数或设备函数）调用，不能由CPU直接调用。 |
3. 执行方式
| 限定符 | 执行方式 |
|---|---|
__global__ | 在GPU上并行执行，由大量线程同时运行。每个线程独立执行核函数中的代码。 |
__device__ | 在GPU上串行执行，由单个线程调用和执行。通常用于实现核函数中的辅助功能。 |
4. 返回值
| 限定符 | 返回值 |
|---|---|
__global__ | 必须返回 void，不能有返回值。计算结果通过修改设备内存或传递指针来实现。 |
__device__ | 可以有返回值，通常用于计算并返回结果。 |
5. 示例对比
__global__ 示例
__global__ 函数是核函数，由CPU调用并在GPU上并行执行：
__global__ void myKernel(int *data, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
data[idx] = data[idx] * 2; // 修改设备内存中的数据
}
}
int main() {
int N = 100;
int *d_data;
cudaMalloc(&d_data, N * sizeof(int)); // 在GPU上分配内存
// 定义网格和线程块
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
// 调用核函数
myKernel<<<blocksPerGrid, threadsPerBlock>>>(d_data, N);
// 将结果从GPU拷贝回CPU
int h_data[N];
cudaMemcpy(h_data, d_data, N * sizeof(int), cudaMemcpyDeviceToHost);
// 释放GPU内存
cudaFree(d_data);
return 0;
}
__device__ 示例
__device__ 函数是设备函数，由核函数或其他设备函数调用：
__device__ int square(int x) {
return x * x; // 计算并返回结果
}
__global__ void myKernel(int *data, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
data[idx] = square(data[idx]); // 调用设备函数
}
}
int main() {
int N = 100;
int *d_data;
cudaMalloc(&d_data, N * sizeof(int)); // 在GPU上分配内存
// 定义网格和线程块
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
// 调用核函数
myKernel<<<blocksPerGrid, threadsPerBlock>>>(d_data, N);
// 将结果从GPU拷贝回CPU
int h_data[N];
cudaMemcpy(h_data, d_data, N * sizeof(int), cudaMemcpyDeviceToHost);
// 释放GPU内存
cudaFree(d_data);
return 0;
}
6. 主要区别总结
| 特性 | __global__ | __device__ |
|---|---|---|
| 用途 | 核函数，GPU程序的入口点 | 设备函数，GPU内部的辅助函数 |
| 调用方式 | 由CPU调用（使用 <<<>>> 语法） | 由GPU上的其他函数调用 |
| 执行方式 | 并行执行，由大量线程同时运行 | 串行执行，由单个线程调用 |
| 返回值 | 必须返回 void | 可以有返回值 |
| 典型应用场景 | 并行计算任务（如矩阵乘法、向量加法等） | 辅助计算（如数学运算、逻辑判断等） |
7. 联合使用
在实际编程中，__global__ 和 __device__ 通常会联合使用。核函数负责并行任务的调度，而设备函数负责实现具体的计算逻辑。例如：
__device__ int add(int a, int b) {
return a + b;
}
__global__ void vectorAdd(int *A, int *B, int *C, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
C[idx] = add(A[idx], B[idx]); // 调用设备函数
}
}
错误处理
所有编程都需要对错误进行处理，早起的编码错误，编译器会帮我们搞定，内存错误也能观察出来，但是有些逻辑错误很难发现，甚至到了上线运行时才会被发现，而且有些厉害的bug复现会很难，不总出现，但是很致命，而且CUDA基本都是异步执行的，当错误出现的时候，不一定是哪一条指令触发的，这一点非常头疼；这时候我们就需要对错误进行防御性处理了，例如代码库头文件里面的这个宏
#define CHECK(call)\
{\
const cudaError_t error=call;\
if(error!=cudaSuccess)\
{\
printf("ERROR: %s:%d,",__FILE__,__LINE__);\
printf("code:%d,reason:%s\n",error,cudaGetErrorString(error));\
exit(1);\
}\
}
就是获得每个函数执行后的返回结果，然后对不成功的信息加以处理，CUDA C 的API每个调用都会返回一个错误代码，这个代码我们就可以好好利用了，当然在release版本中可以去除这部分，但是开发的时候一定要有的。
扫一扫
2136
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
