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

本文用来记录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.pswp.cn/pingmian/82350.shtml

如若内容造成侵权/违法违规/事实不符，请联系多彩编程网进行投诉反馈email:809451989@qq.com，一经查实，立即删除！

相关文章

系统架构中的限流实践：构建多层防护体系（二）

系统架构中的限流实践：构建多层防护体系（二）

系统架构中的限流实践：构建多层防护体系 一、接入层限流：流量拦截第一关二、应用层限流（服务内限流）Java生态方案对比三、分布式限流（跨服务限流）四、数据层限流（数据库/缓存限流）1. 数据库防护策略2. 缓存优化方案五、中间件层限流（消息队列/分布式服务）六、客户端限…
阅读更多...
AI学习笔记二十八:使用ESP32 CAM和YOLOV5实现目标检测

AI学习笔记二十八:使用ESP32 CAM和YOLOV5实现目标检测

若该文为原创文章&#xff0c;转载请注明原文出处。 最近在研究使用APP如何显示ESP32 CAM的摄像头数据&#xff0c;看到有人实现把ESP32 CAM的数据流上传&#xff0c;通过YOLOV5来检测&#xff0c;实现拉流推理&#xff0c;这里复现一下。 一、环境 arduino配置esp32-cam开发环…
阅读更多...
uni-app（5）：Vue3语法基础上

uni-app（5）：Vue3语法基础上

Vue (读音 /vjuː/&#xff0c;类似于 view) 是一套用于构建用户界面的渐进式框架。与其它大型框架不同的是&#xff0c;Vue 被设计为可以自底向上逐层应用。Vue.js 的核心是一个允许采用简洁的模板语法来声明式地将数据渲染进 DOM 的系统&#xff0c;只关注视图层&#xff0c;…
阅读更多...
JAVA：Kafka 存储接口详解与实践样例

JAVA：Kafka 存储接口详解与实践样例

📦 1、简述 Kafka 以其高吞吐、可扩展和高可靠性著称，其强大性能的背后核心在于其高效的存储设计。Kafka 不是传统意义上的队列，而是一个分布式日志系统，其存储模块是核心组成部分。 本文将深入剖析 Kafka 的存储接口实现机制，并结合 Java 示例进行模拟验证。 🧱 2、…
阅读更多...
Docker 使用镜像[SpringBoot之Docker实战系列] - 第537篇

Docker 使用镜像[SpringBoot之Docker实战系列] - 第537篇

历史文章&#xff08;文章累计530&#xff09; 《国内最全的Spring Boot系列之一》 《国内最全的Spring Boot系列之二》 《国内最全的Spring Boot系列之三》 《国内最全的Spring Boot系列之四》 《国内最全的Spring Boot系列之五》 《国内最全的Spring Boot系列之六》 《…
阅读更多...
数据库入门教程：以商品订单系统为例

数据库入门教程：以商品订单系统为例

数据库入门教程&#xff1a;以商品订单系统为例 一、前言 数据库是现代软件开发中不可或缺的基础&#xff0c;掌握数据库的基本概念和操作&#xff0c;是每个开发者的必经之路。本文将以“商品-品牌-客户-订单-订单项”为例&#xff0c;带你快速入门数据库的核心知识和基本操…
阅读更多...
GeoServer样式设置：使用本地图标及分层/分视野显示

GeoServer样式设置：使用本地图标及分层/分视野显示

GeoServer样式设置：使用本地图标及分层/分视野显示 1、本地图标生效设置2、GeoServer添加不同视野的图标点样式1）服务预览效果2）本地图标引用3）不同视野显示不同图标4）标注注记显示空间的点数据，使用图标来表示是非常常见的业务需求，而且由于在不同比例尺下，可能需要设…
阅读更多...
DL00347-基于人工智能YOLOv11的安检X光危险品刀具检测含数据集

DL00347-基于人工智能YOLOv11的安检X光危险品刀具检测含数据集

&#x1f6a8; AI技术革新&#xff0c;提升安检效率与安全性&#xff01;YOLOv11助力X光危险品刀具检测&#xff01; &#x1f4a1; 在安全领域&#xff0c;效率与精准度的要求从未如此迫切。作为科研人员&#xff0c;是否一直在寻找一款可以提升安检准确率、减少人工干预、提…
阅读更多...
测试计划与用例撰写指南

测试计划与用例撰写指南

测试计划与用例撰写指南 一、测试计划&#xff1a;项目测试的 “导航地图”1.1 测试计划的核心目标 1.2 测试计划的关键要素 1.2.1 项目概述 1.2.2 测试策略 1.2.3 资源与进度 1.2.4 风险评估与应对 二、测试用例&#xff1a;测试执行的 “行动指南”2.1 测试用例的设计原则 2…
阅读更多...
微服务的应用案例

微服务的应用案例

从“菜市场”到“智慧超市”&#xff1a;一场微服务的变革之旅 曾经&#xff0c;我们的系统像一个熙熙攘攘的传统菜市场。所有功能模块&#xff08;摊贩&#xff09;都挤在一个巨大的单体应用中。用户请求&#xff08;买菜的顾客&#xff09;一多&#xff0c;整个市场就拥堵不堪…
阅读更多...
Java设计模式之观察者模式：从基础到高级的全面解析

Java设计模式之观察者模式：从基础到高级的全面解析

文章目录 一、观察者模式基础概念1.1 什么是观察者模式？1.2 观察者模式的四大角色1.3 观察者模式类图二、观察者模式实现步骤2.1 基础实现步骤2.2 详细代码实现第一步：定义主题接口第二步：定义观察者接口第三步：创建具体主题第四步：创建具体观察者第五步：客户端使用三、观…
阅读更多...
GATT 服务的核心函数bt_gatt_discover的介绍

GATT 服务的核心函数bt_gatt_discover的介绍

目录 概述 1 GATT 基本概念 1.1 GATT 的介绍 1.2 GATT 的角色 1.3 核心组件 1.4 客户端操作 2 bt_gatt_discover函数的功能和应用 2.1 函数介绍 2.1 发现类型&#xff08;Discover Type&#xff09; 3 典型使用流程 3.1 服务发现示例 3.2 级联发现模式 3.3 按UUID过…
阅读更多...
【更新至2023年】1985-2023年全国及各省就业人数数据（无缺失）

【更新至2023年】1985-2023年全国及各省就业人数数据（无缺失）

1985-2023年全国及各省就业人数数据&#xff08;无缺失&#xff09; 1、时间&#xff1a;1985-2023年 2、来源&#xff1a;Z国统计年鉴、各省年鉴、新中国60年 3、指标&#xff1a;就业人数 4、范围&#xff1a;全国及31省 5、缺失情况&#xff1a;无缺失 6、指标解释&am…
阅读更多...
0基础学习Linux之揭开朦胧一面：环境基础开发工具

0基础学习Linux之揭开朦胧一面：环境基础开发工具

目录 Linux下安装软件的方案&#xff1a; 对于操作系统的理解&#xff1a; 操作系统的生态问题&#xff1a; 什么是好的操作系统&#xff08;os&#xff09;&#xff1a; 重新理解centos VS ubnutu VS kail&#xff1a; 关于yum: 用 yum 安装软件(安装和卸载软件一定要有r…
阅读更多...
YOLO 算法详解：实时目标检测的里程碑

YOLO 算法详解：实时目标检测的里程碑

在计算机视觉领域&#xff0c;目标检测一直是一个关键且热门的研究方向&#xff0c;而 YOLO&#xff08;You Only Look Once&#xff09;算法凭借其出色的实时性和较高的检测精度&#xff0c;成为了目标检测算法中的明星选手。本文将深入探讨 YOLO 算法的原理、发展历程、技术优…
阅读更多...
leetcode98.验证二叉搜索树：递归法中序遍历的递增性验证之道

leetcode98.验证二叉搜索树：递归法中序遍历的递增性验证之道

一、题目深度解析与BST核心性质 题目描述 验证二叉搜索树&#xff08;BST&#xff09;是算法中的经典问题&#xff0c;要求判断给定的二叉树是否满足BST的定义&#xff1a; 左子树中所有节点的值严格小于根节点的值右子树中所有节点的值严格大于根节点的值左右子树本身也必须…
阅读更多...
MathQ-Verify：数学问题验证的五步流水线，为大模型推理筑牢数据基石

MathQ-Verify：数学问题验证的五步流水线，为大模型推理筑牢数据基石

MathQ-Verify&#xff1a;数学问题验证的五步流水线&#xff0c;为大模型推理筑牢数据基石 大语言模型在数学推理领域进展显著&#xff0c;但现有研究多聚焦于生成正确推理路径和答案&#xff0c;却忽视了数学问题本身的有效性。MathQ-Verify&#xff0c;通过五阶段流水线严格…
阅读更多...
八股战神-JVM知识速查

八股战神-JVM知识速查

1.JVM组成 JVM由那些部分组成&#xff0c;运行流程是什么&#xff1f; JVM是Java程序的运行环境 组成部分&#xff1a; 类加载器&#xff1a;加载字节码文件到内存 运行时数据区&#xff1a;包括方法区&#xff0c;堆&#xff0c;栈&#xff0c;程序计数器&#xff0c;本地…
阅读更多...
Maven：在原了解基础上对pom.xml文件进行详细解读

Maven：在原了解基础上对pom.xml文件进行详细解读

一、pom.xml文件 就像项目管理软件 Make 的 MakeFile、Ant 的 build.xml 一样&#xff0c;Maven 项目的核心是 pom.xml。POM( Project Object Model&#xff0c;项目对象模型 ) 定义了项目的基本信息&#xff0c;用于描述项目如何构建&#xff0c;声明项目依赖&#xff0c;等等…
阅读更多...
Spring Cloud项目登录认证从JWT切换到Redis + UUID Token方案

Spring Cloud项目登录认证从JWT切换到Redis + UUID Token方案

背景介绍 在传统的Spring Boot项目中&#xff0c;用户登录认证常见的方案是使用JWT&#xff08;JSON Web Token&#xff09;来实现无状态的身份验证。JWT凭借自包含用户信息、方便前后端分离、性能较好等优势被广泛采用。 然而&#xff0c;在实际项目中&#xff0c;JWT也有一…
阅读更多...
最新文章

Copyright @ 2022~2023