luoganttcc的博客

这段PTX代码实现了CUDA中计算线程坐标(row,col)的核心逻辑。它分为x和y两个方向计算，每个方向4条指令：首先将blockDim、blockIdx、threadIdx值加载到寄存器，然后通过乘加指令(mad.lo.s32)计算全局坐标。x方向计算列号(col=blockIdx.xblockDim.x+threadIdx.x)，y方向计算行号(row=blockIdx.yblockDim.y+threadIdx.y)。这种分段计算方式体现了PTX作为低级中间表示的机械特性，与CUDA的高层抽象形成对

摘要：本文详细讲解了CUDA中grid/block到矩阵索引的映射关系，以8×8矩阵加法为例。配置采用4×4的block和2×2的grid，完整覆盖矩阵。核心公式展示了blockIdx和threadIdx如何组合计算矩阵坐标，并通过区域划分图直观展示映射关系。文章提供了完整的CUDA kernel代码示例，实现C[row][col]=A[row][col]+B[row][col]的并行计算，包含边界检查。最后指出这种映射关系直接影响内存访问效率、warp执行等性能关键因素，是理解GPU并行计算的基础。

本文系统阐述了高性能计算中的索引映射核心公式 global = block × blockSize + local 的数学本质与工程应用。该公式本质是一种多进制展开，将分块坐标（高位）与块内坐标（低位）线性组合为连续地址，其中blockSize作为权重（stride）决定访问步长。文中详细对比了该公式与进制系统、CUDA线程索引、矩阵内存布局的统一性，指出其在高维张量、编译器优化等场景的普适性。特别强调索引展开的连续性对性能的关键影响：合理权重（如stride=1）可实现内存合并访问，而错误权重会导致带宽利

本文分析了CUDA任务在A100 GPU上的切分与执行过程。首先通过示例代码计算总任务规模：24个block，每个block128线程，共3072线程，分为96个warp。重点阐述了block内warp的切分方式，在block(4,8,4)配置下，每个z平面正好对应一个warp（4×8=32线程）。文章详细展示了线性thread id的计算公式及其展开方式，说明warp沿x→y连续扩展，z维负责切分warp。最后指出A100执行时以warp为调度单位，并分析该配置的性能特点：结构清晰但总warp数偏少。理解

摘要： NVIDIA RTX 4090 与 A100 的性能对比需分场景讨论。在单卡、计算密集型任务（如中小模型推理、图像生成）中，4090凭借更多SM单元、更高频率和更强单卡算力（FLOPs）表现更优。然而，在内存/带宽敏感（如大模型推理、长序列）、多卡协同或训练场景下，A100凭借HBM显存、NVLink及高系统吞吐能力显著领先。核心差异在于：4090擅长单点爆发，但受限于显存和通信；A100则强在持续供给与多卡扩展，体现“系统吞吐＞峰值算力”的工程逻辑。

NVIDIA RTX 4090与A100的核心对比显示：4090采用Ada架构，拥有128个SM和16384个CUDA核心，是A100(108SM/6912核心)的2.3倍，理论算力更强。但A100凭借HBM2e显存(1.5-2TB/s带宽)、NVLink和多卡扩展能力，在大规模AI任务中表现更优。关键差异在于4090侧重算力堆叠，而A100实现系统级平衡（算力+带宽+通信），特别适合数据中心场景。最终结论：4090适合计算密集型任务，A100则在大模型训练等场景中更具优势。

本文分析了CUDA中3D网格和块配置的执行模型，重点讨论了warp对齐问题。以一个典型配置dim3 block_size(4,8,4)为例，每个block包含128个线程（4×8×4），正好划分为4个完整的32线程warp。特别指出每个z层（tz）对应一个warp，这种结构完全符合warp=32的要求且没有资源浪费。文章还解释了线程线性化计算方式、warp划分逻辑，并强调虽然这种配置结构规整，但实际性能还需考虑内存访问连续性、共享内存使用等因素。最后提供了一个完整的CUDA示例代码，展示如何打印3D线程布局

本文介绍了在A100 GPU上优化FlashAttention tile size的工程方法。核心结论是：常见配置为BLOCK_M/BLOCK_N=64/128，BLOCK_K=32/64/128，总线程数128或256。这些参数由四个关键资源决定：共享内存容量、寄存器压力、warp数量和load/compute重叠能力。文章提供了A100 SM的关键资源数据：每个SM最多2048线程、64 warps、65536个32-bit寄存器和约160KB共享内存。优化过程需要平衡这些资源限制，计算每个block占

文章摘要： 共享内存（shared memory）本质上是GPU芯片上的SRAM（静态随机存储器），属于高速片上存储。在NVIDIA架构中，它作为SM（流式多处理器）内部的可编程存储区域，具有比HBM（高带宽内存）更快的访问速度。关键区别在于：共享内存是片上SRAM，延迟极低；而HBM是片外存储，带宽高但延迟较大。这种设计使得共享内存特别适合需要频繁数据交换的并行计算任务。

GPU中的warp（32个线程）执行机制解析 摘要： 本文澄清了关于GPU中warp执行的常见误解。一个warp包含32个线程，这些线程会被映射到一组执行单元（core pipeline）并行执行，但并非固定占用32个物理核心。在理想情况下，当SM（流式多处理器）拥有足够计算单元时，32个线程可以同时执行，近似相当于32个核心的并行计算能力。然而实际执行中，warp的线程调度还受限于硬件资源分配和指令流水线等因素，不能简单等同于32个物理核心的固定占用。正确理解warp的执行机制对于优化GPU程序性能至关重

GPU调度常见误区：gridDim不需要是SM数量的整数倍 核心结论： 误区纠正：gridDim（网格维度）不必是SM（流多处理器）数量的整数倍 关键要求：gridDim应远大于SM数量才能充分利用GPU算力 常见误解来源： 直觉认为"1个SM对应1个block"是最优分配 实际GPU采用动态调度机制，每个SM会并行处理多个block 调度机制说明： 现代GPU采用warp调度器，SM会同时处理多个block 需要足够多的block（远大于SM数）才能保持所有SM持续满载 典型的经验值是

1.下载CUDA Toolkit。2.按照命令下载，安装。

【代码】最小 CUDA 示例 cudaMalloc。

link

大致的知道要coalesced access，用vectorized datatype，以及大致的知道 atomics 慢，要避免 bank conflict， 但不是很懂architecture。我的领导就是我见过的最纯粹的体系结构派，碰到问题不是先写代码，而是先建一个excel表格来估算我们最好/最差的情况预期能达到怎样的效率。我觉得计算机专业出身的话，可能学习的路径是反过来的。更新下上手之后对我帮助比较大的一些资料，这个时候对计算机底层不是很了解，做的东西还是浮于算法表面。以上是我自己的学习路径。

【代码】终于彻底解决了conda install速度慢的问题。

【代码】终于彻底解决了conda install速度慢的问题。

推荐几个不错的CUDA入门教程（非广告）godweiyang​​字节跳动 AI Lab NLP算法工程师​关注他1,089 人赞同了该文章❝ 最近因为项目需要，入坑了CUDA，又要开始写很久没碰的C++了。对于CUDA编程以及它所需要的GPU、计算机组成、操作系统等基础知识，我基本上都忘光了，因此也翻了不少教程。这里简单整理一下，给同样有入门需求的同学们参考一下。 ❞官方文档及书籍英文好、时间充裕的同学可以精读官方文档或者著作。NVIDIA CUDA C++ Programming Guide「地址：」 h

代码在git#include "error.cuh"#include <stdio.h>int main(int argc, char *argv[]){ int device_id = 0; if (argc > 1) device_id = atoi(argv[1]); //CHECK(cudaSetDevice(device_id)); cudaDeviceProp prop; //CHECK(cudaGetDeviceProper

代码在gitnvcc add1cpu.cu -o add1cpu./add1cpuTime = 352.801 ms.Time = 222.614 ms.Time = 223.758 ms.Time = 224.543 ms.Time = 223.666 ms.Time = 222.695 ms.Time = 223.809 ms.Time = 222.556 ms.Time = 222.363 ms.Time = 221.945 ms.Time = 222.848 ms.Ti

error.cuh#pragma once#include <stdio.h>#define CHECK(call) \do \{ \ const cudaError_t error_code = call

In this chapter, we show how to check CUDA runtime API functions and CUDA kernels.4.1 A macro function checking CUDA runtime API functionsIn the last chapter, we have learned some CUDA runtime API functions, such as cudaMalloc, cudaFree, and cudaMemcpy..

#include <math.h>#include <stdio.h>const double EPSILON = 1.0e-15;const double a = 1.23;const double b = 2.34;const double c = 3.57;void __global__ add1(const double *x, const double *y, double *z, const int N);void __global__ add2(cons

#include <math.h>#include <stdio.h>#include <iostream>using namespace std;const double EPSILON = 1.0e-15;const double a = 1.23;const double b = 2.34;const double c = 3.57;void __global__ add(const double *x, const double *y, doub

简单CUDA程序的基本框架包含头文件const 或宏定义C++ 函数和CUDA 内核的声明主函数 （）{分配主机和设备内存初始化主机内存中的数据将数据从主机传输到设备启动（调用）内核在设备中进行计算将数据从设备传输到主机释放主机和设备内存}C++ 函数和CUDA 内核的定义...

#include <math.h>#include <stdlib.h>#include <stdio.h>#include <iostream>using namespace std;const double EPSILON = 1.0e-15;const double a = 1.23;const double b = 2.34;const double c = 3.57;void add(const double *x, const

#include <stdio.h>#include <iostream>using namespace std;__global__ void hello_from_gpu(){ const int b = blockIdx.x; const int c = blockIdx.y; const int tx = threadIdx.x; const int ty = threadIdx.y; // cout<<b<

#include <stdio.h>#include <iostream>using namespace std;__global__ void hello_from_gpu(){ const int b = blockIdx.x; const int tx = threadIdx.x; const int ty = threadIdx.y; // cout<<b<<endl; printf("Hello Wor

#include <stdio.h>__global__ void hello_from_gpu(){ const int bid = blockIdx.x; const int tid = threadIdx.x; printf("Hello World from block %d and thread %d!\n", bid, tid);}int main(void){ hello_from_gpu<<<2, 4>>

#include <stdio.h>__global__ void hello_from_gpu(){ printf("Hello World from the GPU!\n");}int main(void){ hello_from_gpu<<<2, 4>>>(); cudaDeviceSynchronize(); return 0;}nvcc hello3.cu -o hello3Hello Worl

Error: unsupported compiler: 9.4.0. Use --override to override this check. Missing recommended library: libGLU.so Missing recommended library: libX11.so Missing recommended library: libXi.so Missing recommended library: libXmu.so Missing recommended libra在

Ubuntu环境下挂载新硬盘Ubuntu环境下挂载新硬盘可以参考这两个教程，在挂载的时候，下面的代码报错sudo mount -t ext4 /dev/sdb /devdata原因是硬盘要挂载在某个文件夹下面mkdir /mnt/devdatasudo mount -t ext4 /dev/sdb /devdata...

dim3 grid(3, 2);dim3 block(5, 3);可以转置一下理解#include <stdio.h>#include <iostream>using namespace std;__global__ void hello_from_gpu(){ const int b = blockIdx.x; const int c = blockIdx.y; const int tx = threadIdx.x; co

#include <stdio.h>#include <iostream>using namespace std;__global__ void hello_from_gpu(){ const int b = blockIdx.x; const int tx = threadIdx.x; const int ty = threadIdx.y; // cout<<b<<endl; printf("Hello Wor

git#include <stdio.h>__global__ void hello_from_gpu(){ const int bid = blockIdx.x; const int tid = threadIdx.x; printf("Hello World from block %d and thread %d!\n", bid, tid);}int main(void){ hello_from_gpu<<<2, 4>

物理概念：streaming processor(sp): 最基本的处理单元。GPU进行并行计算，也就是很多个sp同时做处理。现在SP的术语已经有点弱化了，而是直接使用thread来代替。一个SP对应一个threadWarp：warp是SM调度和执行的基础概念，通常一个SM中的SP(thread)会分成几个warp(也就是SP在SM中是进行分组的，物理上进行的分组)，一般每一个WARP中有32个thread.这个WARP中的32个thread(sp)是一起工作的，执行相同的指令，如果没有这么多t

link❝ 最近因为项目需要，入坑了CUDA，又要开始写很久没碰的C++了。对于CUDA编程以及它所需要的GPU、计算机组成、操作系统等基础知识，我基本上都忘光了，因此也翻了不少教程。这里简单整理一下，给同样有入门需求的同学们参考一下。...

Chapter 3 Basic framework of simple CUDA programs3.1 An example: adding up two arraysWe consider a simple task: adding up two arrays of the same length (same number of elements). We first write a C++ program add.cpp solving this problem. It can be compil

Chapter 1: Introduction to GPU and CUDA1.1 Introduction to GPUGPU means graphics processing unit, which is usually compared to CPU (central processing unit). While a typical CPU has a few relatively fast cores, a typical GPU has hundreds or thousands of