当前位置： 首页 > news >正文

CUDA编程笔记（1）--最简单的核函数

news 2025/6/8 4:13:22

本文用来记录cuda编程的一些笔记以及知识

本笔记运行在windows系统，vs编译器中，cuda版本是12.6

先看一下最基本的代码例子：

#include<iostream>
#include<cstdio>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"//它声明了你在 CUDA kernel 函数中会使用的一些 内置变量
__global__ void kernel() {printf("hello world");
}
int main() {kernel <<<1, 1 >>> ();cudaError_t err = cudaDeviceSynchronize();if (err != cudaSuccess) {std::cerr << "CUDA Error: " << cudaGetErrorString(err) << std::endl;}return 0;
}

下面我会介绍所有出现的东西，以让大家对cuda编程有一个最基本的理解

#include "cuda_runtime.h"

✅ 提供 CUDA 运行时 API（Runtime API）的声明

该头文件包含了调用 CUDA 核函数所需的所有基本功能声明，例如：

🚀 设备管理类函数

  • cudaGetDevice()

  • cudaSetDevice()

  • cudaGetDeviceCount()

🧠 内存管理函数

  • cudaMalloc()：在 GPU 上分配内存

  • cudaFree()：释放 GPU 上的内存

  • cudaMemcpy()：主机和设备之间的数据拷贝

🧮 执行控制函数

  • cudaDeviceSynchronize()：等待所有 GPU 上的任务完成

  • cudaDeviceReset()：复位设备（常用于程序退出前）

  • cudaGetLastError()：获取上一次 CUDA 操作的错误

✅ 支持 kernel 函数的 <<<>>> 调用语法

没有这个头文件，编译器就不会正确识别：

kernel<<<numBlocks, threadsPerBlock>>>(...);

这种语法就是 CUDA 的核函数启动方式，必须在包含 cuda_runtime.h 后才有效。

✅ 加载所有 device API 所需的类型和结构体

比如：

  • cudaError_t：CUDA 错误码的类型

  • cudaMemcpyKind：用于标识拷贝方向（如 cudaMemcpyHostToDevice）

 #include "device_launch_parameters.h"

✅ 提供 CUDA 核函数启动时需要的一些内置变量声明

这个头文件声明了你在 CUDA 核函数（__global__ 函数）中常用的以下 线程索引相关变量：

变量名说明
threadIdx当前线程在 block 中的索引（例如 threadIdx.x）
blockIdx当前 block 在 grid 中的索引
blockDim当前 block 中线程的维度
gridDim当前 grid 的 block 数量

这些变量虽然在 CUDA 编译器中是“内置”的，但如果没有包含适当的头文件，有些编译器（或静态分析工具）会报声明缺失的警告。这个头文件就是为了避免此类问题。

✅ 为什么要显式 include 它？

虽然在大多数情况下 #include <cuda_runtime.h> 已经间接包含了它，但显式写出来有几个好处：

  1. 可读性更好：清晰知道 kernel 中用到哪些内置变量；

  2. 避免 IDE 报错：有些开发工具（如 Visual Studio）如果没有包含这个头文件，在编辑器中会提示这些变量未声明；

  3. 提高代码可移植性：一些平台或旧版本 CUDA SDK 不自动包含它。

核函数

✅ 一、核函数的基本语法

🔹 定义核函数

__global__ void kernel_name(parameter_list) { // GPU 上执行的代码 }

  • __global__：CUDA 的函数修饰符，表示该函数：

    • 在 GPU 上执行

    • 由 CPU 主机代码调用

  • kernel_name：核函数的名称。

  • parameter_list：参数列表，可以传递普通数据指针（如设备内存地址）或标量值。

🔹 调用核函数（使用 <<< >>> 执行配置）

kernel_name<<<numBlocks, threadsPerBlock>>>(parameter_values);

  • numBlocks：grid 中的 block 数量（可以是一维、二维或三维）。

  • threadsPerBlock：每个 block 中的线程数量（同样支持一维、二维或三维）。

  • parameter_values：传递给核函数的参数。

✅ 示例（最常见的一维）：
__global__ void addKernel(int* data) { 
int idx = blockIdx.x * blockDim.x + threadIdx.x; data[idx] += 1; 
} 
int main() {addKernel<<<10, 256>>>(device_ptr); // 共启动 2560 个线程 
}

✅ 二、核函数中常用的内置变量

变量名类型说明
threadIdxdim3当前线程在其所属 block 中的索引
blockIdxdim3当前 block 在 grid 中的索引
blockDimdim3当前 block 中线程的数量
gridDimdim3当前 grid 中 block 的数量

注意：这些变量都是结构体 dim3 类型，支持 .x, .y, .z 维度访问。

✅ 示例：全局索引计算

int globalIndex = blockIdx.x * blockDim.x + threadIdx.x;

🚧 三、核函数的注意事项

1. 核函数只能由 CPU 调用，不能被设备端其他函数调用

  • 核函数必须使用 __global__ 修饰；

  • 不能在 GPU 设备端用另一个核函数调用它；

  • 如果你想在设备端调用，可以使用 __device__ 或 __host__ __device__ 修饰的函数。

2. 核函数返回类型必须是 void

__global__ void foo() {}// ✅ 合法 
__global__ int foo() { return 1; } ///❌ 不合法

3. 不能使用 std::cout，只能使用 printf()

  • 核函数中不支持 C++ 的标准流输出。

  • 使用 #include <cstdio> 和 printf() 进行调试。

4. 核函数中的数组必须放在共享内存或全局内存中

__global__ void kernel() { 
int arr[10]; // ✅ 合法，小数组在寄存器中 // 大数组建议用 `__shared__` 或 `cudaMalloc`}

5. 核函数必须同步（如 cudaDeviceSynchronize()）

  • 如果主机代码紧接着要访问 GPU 数据，必须使用：

cudaDeviceSynchronize();

6. 线程索引越界要自己判断！

  • GPU 不会自动防止访问越界；

  • 通常你要在核函数开头写判断逻辑：

7. 设备函数（__device__）可以在核函数中调用

__device__ int square(int x) { return x * x; }__global__ void kernel(int* out) {int idx = threadIdx.x;out[idx] = square(idx);
}

✅ 总结：核函数语法关键点

元素用法示例
定义方式__global__ void kernel(...)
启动语法kernel<<<blocks, threads>>>(...)
内置索引变量blockIdx, threadIdx, blockDim, gridDim
索引计算int idx = blockIdx.x * blockDim.x + threadIdx.x
输出方式printf()
同步函数cudaDeviceSynchronize()

grid，block，thread

✅ 一、三者之间的关系

Grid（网格）
└── 多个 Block（线程块）
     └── 多个 Thread（线程）

  • Grid：一次 kernel 启动时生成的所有线程块的集合。

  • Block：由多个线程组成的执行单元，GPU 以 block 为基本调度单位。

  • Thread：最基本的并行执行单元。

✅ 二、它们的作用与结构

1. Thread（线程）

  • 执行最小单位，每个线程有自己的局部变量。

  • 每个线程可以根据自己的 ID 来处理不同的数据元素。

CUDA 中线程 ID：

threadIdx.x, threadIdx.y, threadIdx.z

2. Block（线程块）

  • 一个线程块内的线程可以：

    • 使用共享内存（__shared__）进行数据共享

    • 通过 __syncthreads() 进行同步

  • 每个线程块最多可有 1024 个线程（受 GPU 限制）

CUDA 中 block ID：

blockIdx.x, blockIdx.y, blockIdx.z

CUDA 中 block 内线程维度：

blockDim.x, blockDim.y, blockDim.z

3. Grid（网格）

  • Grid 是所有 Block 的集合。

  • 所有的 Block 并行执行。

Grid 维度：

gridDim.x, gridDim.y, gridDim.z

✅ 三、线程索引计算：获取线程全局 ID

一般我们会用下面公式获取线程在整个 Grid 中的全局 ID（以一维为例）：

int global_id = blockIdx.x * blockDim.x + threadIdx.x;

二维情形：

int row = blockIdx.y * blockDim.y + threadIdx.y;

int col = blockIdx.x * blockDim.x + threadIdx.x;

✅ 四、维度示意图（以 2D 举例）

GridDim (2x2)
┌────────────┬────────────┐
│ Block(0,0) │ Block(1,0) │
│ Threads    │ Threads    │
└────────────┴────────────┘
┌────────────┬────────────┐
│ Block(0,1) │ Block(1,1) │
│ Threads    │ Threads    │
└────────────┴────────────┘

每个 Block 里还有：

Threads (e.g., 16x16)
┌──────┐
│ t0   │
│ t1   │
│ ...  │
└──────┘

✅ 五、核函数调用中的 <<<>>> 配置

kernel<<<gridDim, blockDim>>>(...);

  • gridDim：Grid 的尺寸（可以是 int 或 dim3）

  • blockDim：每个 Block 中线程的数量

例如：

dim3 grid(4, 4); // 4x4 个 Block

dim3 block(8, 8); // 每个 Block 8x8 个线程（共 64 个线程）

kernel<<<grid, block>>>(...);

总线程数：4×4×8×8 = 1024

✅ 六、一些常见注意事项

注意点说明
每个 Block 中线程总数限制通常为 1024（可查询设备属性）
线程越界处理你必须手动判断是否越界，否则可能非法访问内存
使用共享内存只能在线程块内部共享，线程间需同步
不同 Block 不共享变量Block 之间不能通信（除非使用 global memory）
线程的层级结构是固定的Grid → Block → Thread 是一层层嵌套的

✅ 七、小结表格

层级名称定义方式标识符（内置变量）
1Grid<<<gridDim, ...>>>gridDim, blockIdx
2Block<<<..., blockDim>>>blockDim, threadIdx
3Thread在核函数中自动生成threadIdx

 在示例中，<<<1,1>>>表示我们只启动了一个线程，因此hello world只会输出一次；我们将代码改为如下：

#include<iostream>
#include<cstdio>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"//它声明了你在 CUDA kernel 函数中会使用的一些 内置变量
__global__ void kernel() {/*int current_thread_id = blockIdx.x * blockDim.x + threadIdx.x;printf("current thread id is :%d\n", current_thread_id);*/printf("hello world\n");
}
int main() {kernel <<<4, 8 >>> ();cudaError_t err = cudaDeviceSynchronize();if (err != cudaSuccess) {std::cerr << "CUDA Error: " << cudaGetErrorString(err) << std::endl;}return 0;
}

grid中有4个block，每个block拥有8个thread，所以总共有32个线程，也就会输出32个hello world。

计算线程的唯一id

线程在线性空间中的全局编号（global thread id）

int global_id =threadIdx.x+ threadIdx.y * blockDim.x+ threadIdx.z * blockDim.x * blockDim.y+ blockIdx.x * blockDim.x * blockDim.y * blockDim.z+ blockIdx.y * gridDim.x * blockDim.x * blockDim.y * blockDim.z+ blockIdx.z * gridDim.x * gridDim.y * blockDim.x * blockDim.y * blockDim.z;

分解解释：

线程在线程块内的编号（local thread offset）：

local_id = threadIdx.x+ threadIdx.y * blockDim.x+ threadIdx.z * blockDim.x * blockDim.y;

线程块的编号（block offset）： 

block_id = blockIdx.x+ blockIdx.y * gridDim.x+ blockIdx.z * gridDim.x * gridDim.y;

每个 Block 中的线程数量：

threads_per_block = blockDim.x * blockDim.y * blockDim.z;

 所以也可以写成：

global_id = block_id * threads_per_block + local_id;

如果你处理的是三维数组索引

除了线性 global_id，你还可以直接求出每个线程在全局的 (x, y, z) 坐标：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int z = blockIdx.z * blockDim.z + threadIdx.z;

类比：想象一组三维数组任务

比如你现在有个数据张量：float data[64][64][64]，你希望：

  • 每个线程负责处理这个数组的一个元素，比如 data[z][y][x]。

  • CUDA 会把这些任务分给多个线程，你就用上面那行代码求出：每个线程该处理哪个 x,y,z 的位置。

具体例子

假设你启动 CUDA kernel去处理一个[8][8][8]的数组（当然也可以是一维的）

kernel<<<dim3(2,2,2), dim3(4,4,4)>>>();

这表示：

  • Grid 有 2x2x2=8 个 Block

  • 每个 Block 有 4x4x4=64 个线程

  • 所以总线程数 = 8 × 64 = 512

那么：

  • 第一个 Block（blockIdx=(0,0,0)）内的线程 threadIdx=(0,0,0) 处理的是 (x=0, y=0, z=0)

  • 第二个 Block（blockIdx=(1,0,0)）内的 threadIdx=(0,0,0) 处理的是 (x=4, y=0, z=0)（x 方向偏移了一个 Block）

  • 第三个 Block（blockIdx=(0,1,0)）内的 threadIdx=(0,0,0) 处理的是 (x=0, y=4, z=0)（y 方向偏移了一个 Block）

  • 第四个 Block（blockIdx=(1,1,0)）内的 threadIdx=(1,2,3) 处理的是 (x=5, y=6, z=3)

cudaDeviceSynchronize的作用

cudaDeviceSynchronize() 是 CUDA 编程中的一个重要函数，它的主要作用是：

阻塞当前 CPU 线程，直到 GPU 上之前所有的任务全部完成。

详细来说：

  • CUDA 的操作（如内核函数启动、内存拷贝等）是异步执行的，调用它们时 CPU 不会等待 GPU 完成任务就直接继续执行后面的代码。

  • cudaDeviceSynchronize() 会让 CPU 阻塞，直到 GPU 上所有之前发起的任务（比如内核执行、数据传输等）都执行完毕，确保 GPU 任务完成后再继续执行 CPU 代码。

  • 它常用于调试、性能测量或者需要确保 GPU 计算结果已经完成才能进行下一步处理的场景。

http://www.xdnf.cn/news/658315.html

相关文章：

  • 大模型RL方向面试题90道
  • Filter和Interceptor详解（一文了解执行阶段及其流程）
  • CVE-2024-36467 Zabbix权限提升
  • java枚举和mybaits-plus结合实现映射输出和存储
  • VScode怎么运行一个c语言程序
  • ChatGPT与认知科学：人机协同的未来图景
  • STM32 IIC总线死锁问题总结
  • 洛谷——P3372 【模板】线段树 1
  • webpack吐环境分析
  • 为什么使用ollama运行的模型不用gpu也可以使用
  • [攻防世界] easyphp writeup
  • Graph Neural Network(GNN)
  • 如何通过全流量溯源分析系统实现高效的网络质量监控
  • JavaSE核心知识点04工具04-02（IDEA）
  • 关于（stream）流
  • MySQL的基础操作
  • 内网搭建NTS服务器
  • 网络安全之Web渗透加解密
  • 原子操作（Atomic Operations）在SOC中的应用场景
  • 【R语言编程绘图-函数篇】
  • Sparse VideoGen开源：完全无损，视频生成速度加速两倍，支持Wan 2.1、HunyuanVideo等
  • DAY12打卡 启发式算法
  • 基于yjs实现协同编辑页面
  • 学习黑客Metasploit 框架的原理
  • 端午假期 · 粽享欢乐
  • 开源Vue表单设计器 FcDesigner 组件提供的方法详解
  • 《1.1_4计算机网络的分类|精讲篇|附X-mind思维导图》
  • deepseek告诉您http与https有何区别？
  • CQF预备知识：一、微积分 -- 1.4.6 莱布尼茨法则详解
  • Mysql在SQL层面的优化