CUDA学习笔记
CUDA基础概念
CUDA 内置变量
1. 线程和块索引变量
这些变量用于确定当前线程在网格（grid）和块（block）中的位置：
| 变量名 | 数据类型 | 说明 |
|---|---|---|
threadIdx |
dim3 |
当前线程在 block 内的索引（x, y, z 分量） |
blockIdx |
dim3 |
当前 block 在 grid 内的索引（x, y, z 分量） |
blockDim |
dim3 |
block 的维度（每个 block 的线程数，x, y, z 分量） |
gridDim |
dim3 |
grid 的维度（每个 grid 的 block 数，x, y, z 分量） |
2. 计算能力相关的常量
这些变量描述了 GPU 的硬件特性：
| 变量名 | 数据类型 | 说明 |
|---|---|---|
warpSize |
int |
一个 warp 的线程数（通常是 32） |
__CUDA_ARCH__ |
int |
编译目标的计算能力（如 700 表示 SM 7.0） |
3. 内存和同步变量
这些变量用于控制内存访问和线程同步：
| 变量名 | 说明 |
|---|---|
__shared__ |
定义共享内存（每个 block 共享） |
__constant__ |
定义常量内存（只读，全局可见） |
__device__ |
定义全局设备变量（所有线程可见） |
__syncthreads() |
块内线程同步（等待所有线程执行到此处） |
4. 数学和特殊函数
CUDA 提供了一些内置数学函数（如 sin, exp, atomicAdd）：
| 类别 | 示例函数 |
|---|---|
| 数学运算 | sin, cos, exp, log, sqrt |
| 原子操作 | atomicAdd, atomicSub, atomicMax |
| 类型转换 | __float2half, __int2float_rn |
5. 其他特殊变量
| 变量名 | 说明 |
|---|---|
clock() |
获取 GPU 时钟周期（用于计时） |
nanosleep() |
线程休眠（ns 级延迟） |
本博客所有文章除特别声明外，均采用 CC BY-NC-SA 4.0 许可协议。转载请注明来源 一只大笨熊！
相关推荐
2025-06-19
常见手撕算子——一维数组的softmax
SoftMax Softmax 的 CPU 和 CUDA 写法均是高频考察。面试时有可能会让任选一种写法进行书写，此时自己可以先写 CPU(C++、Python) 版本，然后再写 CUDA 版本。 Softmax公式如下：softmax(xi)=exi∑jexjsoftmax(x_i) = \frac{e^{x_i}}{\sum_j e^{x_j}} softmax(xi)=∑jexjexi 一般为了避免溢出，需要减去最大值，所以通常采用下面这个公式：softmax(xi)=exi−max(x)∑jexj−max(x)softmax(x_i) = \frac{e^{x_i - max(x)}}{\sum_j e^{x_j - max(x)}} softmax(xi)=∑jexj−max(x)exi−max(x) 1. CPU(C++、Python) 版本 1234567891011void softmax(float* input, float* output, int N){ float max_value =...
2025-06-19
常见手撕算子-elementwise
elementwise elementwise 是最简单的一类算子，其指的是对数据进行逐元素操作，例如将两个等长的数组对应元素相加（add）。另外在深度学习中，激活函数会对输入数据的每个元素求对应激活值，故激活函数也算在 elementwise 范围内。 add 1234567891011121314151617181920212223242526272829303132333435363738394041// 1. 向上取整#define CEIL(a, b) ((a + b - 1) / (b))// 2. FLOAT4，用于向量化访存，以下两种都可以// c写法#define FLOAT4(value) *(float4*)(&(value))// c++写法#define FLOAT4(value) (reinterpret_cast<float4*>(&(value))[0])//naive版int block_size = 1024;int grid_size = CEIL(N,...
2025-06-19
常见手撕算子-reduce
Reduce 算子是指通过对数组中的每个元素进行操作，得到一个输出值的过程。常见的操作包括求和（sum）、取最大值（max）、取最小值（min）等。在 CUDA 中，优化 Reduce 算子可以显著提高计算效率。 1. naive实现 1234567//累加__global__ void reduce1(float* d_in, float* d_out, int N) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx < N) { atomicAdd(d_out, d_in[idx]); }} 2. share mem + 折半规约 12345678910111213141516171819202122232425262728__global__ void reduce2(float* d_in, float* d_out, int N) { __shared__ float sdata[BLOCK_SIZE]; ...
2025-06-19
常见手撕算子——sgemm（单精度矩阵乘法）
1. cpu: 矩阵乘法 123456789101112131415void matrixMultiply(const float* A, const float* B, float* C, int m, int p, int n) { // A is m x p, B is p x n, C is m x n for (int i = 0; i < m; ++i) { for (int j = 0; j < n; ++j) { float sum = 0.0; for (int k = 0; k < p; ++k) { // A[i][k] -> A[i * p + k] // B[k][j] -> B[k * n + j] sum += A[i * p + k] * B[k * n + j]; } ...
2025-06-19
常见手撕算子——transformer的softmax_matrix
1.cpu: 计算每行的softmax 12345678910111213141516void softmax_row(float* input, float* output, int M, int N) { for (int row = 0; row < M; row++) { // 第row行 float* input_tmp = input + row * N; float* output_tmp = output + row * N; float max_val = *(std::max_element(input_tmp, input_tmp + N)); // 计算输入数组的最大值 float sum = 0; for (int i = 0; i < N; i++) { output_tmp[i] = std::exp(input_tmp[i] - max_val); //...
2025-06-19
常见手撕算子-transpose
naive版本 1234567__global__ void transpose_v0(float* input, float* output, int M, int N){ int row = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; if(row < M && col < N){ output[col * M + row] = input[row * N + col]; }} 优化版本1：shared memory 思路： 先将数据从global memory拷贝到shared memory中 通过shared memory进行转置 通过shared memory将数据拷贝到global memory中 1234567891011121314151617template <int TILE_SIZE>__global__...
评论