cuda编程入门
CUDA 入门编程示例:
本文将通过几个简单的 CUDA 编程示例，带你初步了解 CUDA 的基本概念和编程流程。这些示例都附带了详细的解释，希望能帮助你快速入门 CUDA 开发。
1. 第一个 CUDA 核函数：Hello GPU!
让我们从一个最简单的 CUDA 程序开始，这个程序将在 GPU 上打印一条消息。
#include <stdio.h>
// CPU函数：在主机端打印问候消息
void helloCPU()
{
printf("Hello from the CPU.\n");
}
// GPU核函数：在设备端打印问候消息
// __global__ 修饰符表示这是一个可以从主机调用并在GPU上执行的核函数
__global__ void helloGPU()
{
printf("Hello from the GPU!\n");
}
int main()
{
// 调用CPU函数
helloCPU();
// 调用GPU核函数
// <<<1, 1>>> 表示使用1个线程块，每个块中使用1个线程
helloGPU<<<1, 1>>>();
// 等待GPU核函数执行完成
// 这是必需的，因为GPU核函数的执行是异步的
cudaDeviceSynchronize();
return 0;
}
核心概念解析：
__global__修饰符： 这是定义 CUDA 核函数的关键。__global__表明该函数将在 GPU 设备上执行，并且可以从 CPU 主机端调用。任何想要在 GPU 上运行的代码都需要被声明为核函数。- 核函数调用
<<<1, 1>>>： 这部分语法用于启动核函数。<<<grid, block>>>定义了核函数在 GPU 上执行的线程网格（grid）和线程块（block）的维度。在这个例子中，<<<1, 1>>>表示我们启动了一个包含一个线程块的网格，而这个线程块中只包含一个线程。这是最简单的核函数启动方式。 cudaDeviceSynchronize()： GPU 的执行是异步的，这意味着当 CPU 调用核函数后，它不会等待 GPU 执行完成就继续执行后续的代码。cudaDeviceSynchronize()函数的作用是让 CPU 阻塞，直到所有已启动的 GPU 核函数都执行完毕。在简单的示例中，为了确保我们能看到 GPU 的输出，通常需要在核函数调用后加上这个同步操作。
运行这个程序：
你需要使用 NVIDIA CUDA 编译器 nvcc 来编译这个 .cu 文件。通常的编译命令如下：
nvcc hello.cu -o hello
./hello
你应该会看到如下输出：
Hello from the CPU.
Hello from the GPU!
这表明你的第一个 CUDA 核函数已经在 GPU 上成功运行了！
2. 获取 CUDA 设备信息
在开始更复杂的 CUDA 编程之前，了解你的 GPU 设备的一些基本属性是非常有用的。下面的代码展示了如何获取 CUDA 设备的信息。
#include <stdio.h>
int main()
{
// 获取当前CUDA设备的ID
int deviceId;
cudaGetDevice(&deviceId);
// 创建设备属性结构体并获取设备属性
cudaDeviceProp props;
cudaGetDeviceProperties(&props, deviceId);
// 基本设备信息
int computeCapabilityMajor = props.major;
int computeCapabilityMinor = props.minor;
int multiProcessorCount = props.multiProcessorCount;
int warpSize = props.warpSize;
int maxThreadsPerBlock = props.maxThreadsPerBlock;
int maxThreadsPerMultiProcessor = props.maxThreadsPerMultiProcessor;
// 计算Warp相关信息
int warpsPerSM = maxThreadsPerMultiProcessor / warpSize;
int totalWarps = warpsPerSM * multiProcessorCount;
// 获取维度限制
int maxThreadsDim[3] = {props.maxThreadsDim[0],
props.maxThreadsDim[1],
props.maxThreadsDim[2]};
int maxGridSize[3] = {props.maxGridSize[0],
props.maxGridSize[1],
props.maxGridSize[2]};
// 获取内存信息
size_t totalGlobalMem = props.totalGlobalMem;
size_t sharedMemPerBlock = props.sharedMemPerBlock;
// 打印完整设备信息
printf("\n===== CUDA 设备属性 =====\n");
printf("\n基本信息:\n");
printf("设备名称: %s\n", props.name);
printf("设备ID: %d\n", deviceId);
printf("计算能力: %d.%d\n", computeCapabilityMajor, computeCapabilityMinor);
printf("\n处理器信息:\n");
printf("SM数量: %d\n", multiProcessorCount);
printf("每个SM的最大线程数: %d\n", maxThreadsPerMultiProcessor);
printf("每个块的最大线程数: %d\n", maxThreadsPerBlock);
printf("\nWarp信息:\n");
printf("Warp大小: %d\n", warpSize);
printf("每个SM的Warp数量: %d\n", warpsPerSM);
printf("总Warp数量: %d\n", totalWarps);
printf("\n维度限制:\n");
printf("最大线程维度: (%d, %d, %d)\n",
maxThreadsDim[0], maxThreadsDim[1], maxThreadsDim[2]);
printf("最大网格维度: (%d, %d, %d)\n",
maxGridSize[0], maxGridSize[1], maxGridSize[2]);
printf("\n内存信息:\n");
printf("全局内存: %.2f GB\n", (float)totalGlobalMem / (1024*1024*1024));
printf("每块共享内存: %zu KB\n", sharedMemPerBlock / 1024);
printf("\n理论性能:\n");
printf("最大并发线程数: %d\n", multiProcessorCount * maxThreadsPerMultiProcessor);
printf("最大并发Warp数: %d\n", totalWarps);
return 0;
}
核心概念解析：
cudaGetDevice(&deviceId)： 这个函数用于获取当前 CUDA 设备的 ID。通常，如果你的系统只有一个 NVIDIA GPU，那么它的 ID 就是 0。cudaDeviceProp props;和cudaGetDeviceProperties(&props, deviceId)：cudaDeviceProp是一个结构体，用于存储设备的各种属性。cudaGetDeviceProperties()函数用于填充这个结构体，你需要提供设备 ID 和指向cudaDeviceProp结构体的指针。- 设备属性： 代码中打印了许多重要的设备属性，例如：
- 计算能力 (Compute Capability)： 表示 GPU 的硬件架构级别，决定了 GPU 支持的 CUDA 特性和指令集。
- SM 数量 (multiProcessorCount)： Streaming Multiprocessor 的数量，每个 SM 包含多个 CUDA 核心。
- Warp 大小 (warpSize)： CUDA 执行的基本调度单位，通常是 32 个线程。
- 每个块的最大线程数 (maxThreadsPerBlock)： 每个线程块中可以容纳的最大线程数量。
- 全局内存 (totalGlobalMem)： GPU 上可用的总显存大小。
- 共享内存 (sharedMemPerBlock)： 每个线程块可用的共享内存大小，共享内存比全局内存快得多，是优化的关键。
- 最大线程维度 (maxThreadsDim) 和最大网格维度 (maxGridSize)： 定义了线程块和线程网格在三个维度上的最大尺寸。
运行这个程序可以帮助你了解你的 GPU 的性能和限制，这对于后续的 CUDA 程序设计和优化至关重要。
3. 利用全局线程索引防止越界
在实际的并行计算中，我们需要将计算任务分配给大量的 GPU 线程。正确地为每个线程分配数据元素至关重要，以避免越界访问等问题。
#include <stdio.h>
/*
* Initialize array values on the host.
*/
// 在主机端初始化数组：将数组元素设置为其索引值
void init(int *a, int N)
{
int i;
for (i = 0; i < N; ++i)
{
a[i] = i; // 每个元素的值等于其索引
}
}
// GPU核函数：并行将数组中的每个元素翻倍
__global__
void doubleElements(int *a, int N)
{
int i;
// 计算全局线程索引
i = blockIdx.x * blockDim.x + threadIdx.x;
// 确保不越界
if (i < N)
{
a[i] *= 2; // 将元素值翻倍
}
}
/*
* Check all elements have been doubled on the host.
*/
// 在主机端验证所有元素是否都已正确翻倍
bool checkElementsAreDoubled(int *a, int N)
{
int i;
for (i = 0; i < N; ++i)
{
// 检查每个元素是否等于其索引值的两倍
if (a[i] != i*2) return false;
}
return true;
}
int main()
{
// 设置数组大小
int N = 100;
int *a;
size_t size = N * sizeof(int);
// 分配统一内存，使CPU和GPU都能访问
cudaMallocManaged(&a, size);
// 初始化数组
init(a, N);
// 设置CUDA核函数的执行配置
size_t threads_per_block = 10; // 每个块10个线程
size_t number_of_blocks = 10; // 总共10个块
// 启动核函数，并行处理数组
doubleElements<<<number_of_blocks, threads_per_block>>>(a, N);
// 等待GPU操作完成
cudaDeviceSynchronize();
// 验证结果
bool areDoubled = checkElementsAreDoubled(a, N);
printf("All elements were doubled? %s\n", areDoubled ? "TRUE" : "FALSE");
// 释放统一内存
cudaFree(a);
}
核心概念解析：
- 线程块和线程索引： CUDA 将执行的线程组织成线程块（block），而线程块又组织成线程网格（grid）。在核函数中，每个线程都有其唯一的索引：
blockIdx.x：当前线程块在网格中的 x 维度索引。blockDim.x：每个线程块在 x 维度上的线程数量。threadIdx.x：当前线程在线程块中的 x 维度索引。
- 计算全局线程索引： 通过
int i = blockIdx.x * blockDim.x + threadIdx.x;这行代码，我们可以计算出当前线程在整个网格中的全局唯一索引。这使得每个线程可以负责处理数组中特定的元素。 - 越界检查： 由于线程的数量可能大于数组的大小，因此在访问数组元素之前，进行越界检查 (
if (i < N)) 非常重要，以防止程序崩溃或产生未定义行为。 - 统一内存 (
cudaMallocManaged)： 在这个例子中，我们使用了cudaMallocManaged来分配内存。统一内存允许 CPU 和 GPU 直接访问同一块内存，简化了数据在主机和设备之间的传输。
这个示例展示了如何使用线程块和线程索引来并行地处理数组中的每个元素，并确保不会发生越界访问。
4. initWith 函数和 CUDA 错误处理示例
良好的编程实践包括及时的错误处理。CUDA 提供了一些机制来检测和处理 GPU 操作中可能出现的错误。
#include <stdio.h>
// CUDA错误处理宏，用于包装CUDA函数调用
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
// CUDA错误处理辅助函数
inline void gpuAssert(cudaError_t code, const char *file, int line)
{
if (code != cudaSuccess)
{
fprintf(stderr, "GPU错误: %s\n文件: %s\n行号: %d\n",
cudaGetErrorString(code), file, line);
exit(code);
}
}
// CPU函数：用指定值初始化数组
void initWith(float num, float *a, int N)
{
for(int i = 0; i < N; ++i)
{
a[i] = num;
}
}
// GPU核函数：并行执行向量加法
// 使用网格跨度循环处理大数组
__global__
void addVectorsInto(float *result, float *a, float *b, int N)
{
// 计算当前线程的全局索引
int index = blockIdx.x * blockDim.x + threadIdx.x;
// 计算网格总大小（总线程数）
int stride = blockDim.x * gridDim.x;
// 使用网格跨度循环，每个线程处理多个元素
for(int i = index; i < N; i += stride)
{
result[i] = a[i] + b[i];
}
}
// CPU函数：验证结果是否正确
void checkElementsAre(float target, float *array, int N)
{
for(int i = 0; i < N; i++)
{
if(array[i] != target)
{
printf("FAIL: array[%d] - %0.0f does not equal %0.0f\n", i, array[i], target);
exit(1);
}
}
printf("SUCCESS! All values added correctly.\n");
}
int main()
{
// 设置数组大小：2^21 = 2,097,152个元素
const int N = 2<<20;
size_t size = N * sizeof(float);
// 声明三个浮点数数组指针
float *a;
float *b;
float *c;
// 分配统一内存，使CPU和GPU都能访问
gpuErrchk(cudaMallocManaged(&a, size));
gpuErrchk(cudaMallocManaged(&b, size));
gpuErrchk(cudaMallocManaged(&c, size));
// 初始化数组：a全为3，b全为4，c全为0
initWith(3, a, N);
initWith(4, b, N);
initWith(0, c, N);
// 设置CUDA核函数执行配置
int threadsPerBlock = 256;
// 计算需要的块数，确保能处理所有元素
int numberOfBlocks = (N + threadsPerBlock - 1) / threadsPerBlock;
// 启动核函数进行向量加法
addVectorsInto<<<numberOfBlocks, threadsPerBlock>>>(c, a, b, N);
// 检查核函数启动错误
gpuErrchk(cudaPeekAtLastError());
// 等待GPU完成并检查运行时错误
gpuErrchk(cudaDeviceSynchronize());
// 验证结果：所有元素应该等于7（3+4）
checkElementsAre(7, c, N);
// 释放统一内存
gpuErrchk(cudaFree(a));
gpuErrchk(cudaFree(b));
gpuErrchk(cudaFree(c));
return 0;
}
核心概念解析：
initWith函数： 这是一个简单的 CPU 函数，用于将数组的所有元素初始化为指定的数值。在 CUDA 编程中，我们经常需要在主机端初始化数据，然后将其传输到 GPU 进行处理。- CUDA 错误处理：
gpuErrchk(ans)宏： 这是一个自定义的宏，用于简化 CUDA 函数调用的错误检查。它接受一个 CUDA 函数的返回值（cudaError_t类型），并将其传递给gpuAssert函数进行处理。gpuAssert函数： 这个内联函数检查传入的cudaError_t代码。如果代码不是cudaSuccess，则表示发生了错误。函数会打印错误信息（包括错误字符串、文件名和行号）并终止程序。cudaPeekAtLastError()： 用于检查最近一次 CUDA 操作是否发生了错误。cudaDeviceSynchronize()： 除了等待 GPU 完成，它还会返回设备上发生的任何错误。
- 网格跨度循环： 在
addVectorsInto核函数中，我们使用了网格跨度循环来处理可能比总线程数更多的元素。每个线程计算其全局索引index和步长stride（等于总线程数）。然后，线程以stride为步长遍历数组，处理多个元素。这是一种常用的处理大型数据集的技术。
通过使用错误处理机制，我们可以更容易地发现和修复 CUDA 程序中的问题。
5. CUDA 网格（Grid）示例
最后一个示例进一步展示了如何利用 CUDA 的网格和线程块结构来并行处理数据。
#include <stdio.h>
// 在CPU上初始化数组：将每个元素设置为其索引值
void init(int *a, int N)
{
int i;
for (i = 0; i < N; ++i)
{
a[i] = i;
}
}
// GPU核函数：使用网格跨度循环将数组元素翻倍
__global__ void doubleElements(int *a, int N)
{
// 计算总线程数（网格大小）= 每块线程数 × 块数
// 在本例中是 256 × 32 = 8192
int gridSize = blockDim.x * gridDim.x;
// 计算当前线程的全局ID
// 用于确定每个线程的起始处理位置
int threadId = blockIdx.x * blockDim.x + threadIdx.x;
// 网格跨度循环：让每个线程处理多个元素
// i 初始值为线程ID，每次增加网格大小
// 这样可以让每个线程处理多个间隔为gridSize的元素
for (int i = threadId; i < N; i += gridSize)
{
a[i] *= 2;
}
}
// 在CPU上验证数组元素是否都已正确翻倍
bool checkElementsAreDoubled(int *a, int N)
{
int i;
for (i = 0; i < N; ++i)
{
if (a[i] != i*2) return false;
}
return true;
}
int main()
{
// 设置数组大小为10000
// 这个大小大于网格总线程数(8192)
int N = 10000;
int *a;
// 计算所需内存大小
size_t size = N * sizeof(int);
// 分配统一内存，使CPU和GPU都能访问
cudaMallocManaged(&a, size);
// 初始化数组
init(a, N);
// 设置CUDA执行配置
// 总线程数 = 256 × 32 = 8192
size_t threads_per_block = 256; // 每块256个线程
size_t number_of_blocks = 32; // 总共32个块
// 启动核函数
// 使用网格跨度循环处理10000个元素
doubleElements<<<number_of_blocks, threads_per_block>>>(a, N);
// 等待GPU完成
cudaDeviceSynchronize();
// 验证结果
bool areDoubled = checkElementsAreDoubled(a, N);
printf("All elements were doubled? %s\n", areDoubled ? "TRUE" : "FALSE");
// 释放内存
cudaFree(a);
}
核心概念解析：
- 网格大小和线程块大小： 在
main函数中，我们设置了每个线程块包含 256 个线程 (threads_per_block)，并启动了 32 个线程块 (number_of_blocks)。这意味着总共有 256 * 32 = 8192 个 CUDA 线程被创建来执行doubleElements核函数。 - 网格跨度循环的应用： 尽管我们只有 8192 个线程，但数组的大小
N是 10000。为了处理所有元素，核函数内部使用了网格跨度循环。每个线程首先计算其全局 ID (threadId)，然后以整个网格的大小 (gridSize) 为步长遍历数组。这样，每个线程都会负责处理多个相隔gridSize的元素，从而确保所有 10000 个元素都被处理到。
这个示例清晰地展示了如何有效地利用 CUDA 的线程组织结构和网格跨度循环来处理比可用线程总数更多的数据。
总结
通过以上五个简单的 CUDA 编程示例，我们初步了解了 CUDA 核函数的定义和启动、如何获取设备信息、如何使用线程索引防止越界、如何进行基本的错误处理以及如何利用网格跨度循环处理大型数据集。
浙公网安备 33010602011771号