CUDA学习笔记
本博客所有文章除特别声明外，均采用 CC BY-NC-SA 4.0 许可协议。转载请注明来源 一只大笨熊！
相关推荐
2026-08-08
常见手撕算子——一维数组的softmax
SoftMax Softmax 的 CPU 和 CUDA 写法均是高频考察。面试时有可能会让任选一种写法进行书写，此时自己可以先写 CPU(C++、Python) 版本，然后再写 CUDA 版本。 Softmax公式如下： $$softmax(x_i) = \frac{e^{x_i}}{\sum_j e^{x_j}}$$ 一般为了避免溢出，需要减去最大值，所以通常采用下面这个公式： $$softmax(x_i) = \frac{e^{x_i - max(x)}}{\sum_j e^{x_j - max(x)}}$$ 1. CPU(C++、Python) 版本1234567891011void softmax(float* input, float* output, int N){ float max_value = *std::max_element(input, input + N); float sum = 0; for(int i = 0; i < N; i++){ output[i] =...
2026-08-08
常见手撕算子-elementwise
elementwise elementwise 是最简单的一类算子，其指的是对数据进行逐元素操作，例如将两个等长的数组对应元素相加（add）。另外在深度学习中，激活函数会对输入数据的每个元素求对应激活值，故激活函数也算在 elementwise 范围内。 add1234567891011121314151617181920212223242526272829303132333435363738394041// 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,...
2026-08-08
常见手撕算子-reduce
Reduce 算子是指通过对数组中的每个元素进行操作，得到一个输出值的过程。常见的操作包括求和（sum）、取最大值（max）、取最小值（min）等。在 CUDA 中，优化 Reduce 算子可以显著提高计算效率。 1. naive实现1234567//累加__global__ void reduce_v0(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. 使用warp级并行进行数组归约12345678910111213141516171819202122232425262728293031323334353637383940414243444546#include <cuda_runtime.h>#include...
2026-08-08
常见手撕算子——sgemm（单精度矩阵乘法）
1. cpu: 矩阵乘法1234567891011121314151617181920212223242526272829// 二维矩阵void 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) { sum += A[i][k] * B[k][j]; } C[i][j] = sum; } }}// 二维矩阵展开成一维void...
2026-08-08
常见手撕算子——transformer的softmax_matrix
1.cpu: 计算每行的softmax12345678910111213141516void 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); //...
2026-08-08
常见手撕算子-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__...
评论
