1. 什么是算子一个算子Operator / Op就是一个基本计算节点。例如PyTorch 代码c torch.matmul(a, b)从用户角度看是两个矩阵相乘但是 PyTorch 内部会变成MatMul Operator → CUDA Kernel → GPU 执行这里MatMul 是一个算子CUDA Kernel 是真正跑在 GPU 上的代码2. CUDA 算子和 CUDA Kernel 的关系Kernel 是 CUDA 中最底层的 GPU 函数。例如__global__voidadd_kernel(float*a,float*b,float*c){intithreadIdx.x;c[i]a[i]b[i];}调用方式add_kernel...();上述代码详解1.__global__是什么__global__voidadd_kernel()是 CUDA 特有的关键字表示这个函数是一个 CUDA Kernel可以由 CPU 调用但是运行在 GPU 上。2. threadIdx.xintithreadIdx.x;这是 CUDA 最重要的概念之一。GPU 不是一个 CPU 核心执行 for 循环而是很多线程同时运行。例如启动add_kernel1,5();意思是1 个 block5 个线程。GPU 创建Thread 0 Thread 1 Thread 2 Thread 3 Thread 4每个线程都有自己的threadIdx.x于是线程threadIdx.x线程 00线程 11线程 22线程 33线程 44所以线程 0i0线程 1i1线程 2i23. 核心计算c[i]a[i]b[i];由于每个线程有不同的 i线程 0c[0]a[0]b[0]计算11011线程 1c[1]a[1]b[1]计算22022线程 2c[2]a[2]b[2]计算33033最终GPU 线程 Thread 0 c[0]a[0]b[0] Thread 1 c[1]a[1]b[1] Thread 2 c[2]a[2]b[2] Thread3: c[3]a[3]b[3] Thread4: c[4]a[4]b[4]多线程工作原理当前代码intithreadIdx.x;只能处理一个 block 里的线程。比如add_kernel1,256();最多 256 个元素如果数组 1000000 个元素怎么办需要block grid的 CUDA 三层结构Grid | | --- Block 0 | | | Thread0 | Thread1 | --- Block 1 | | | Thread0 | Thread1 | --- Block 2所以真实的 CUDA 代码通常是intiblockIdx.x*blockDim.xthreadIdx.x;也就是计算全局线程编号线程编号 block编号 × 每个block线程数量 线程在线程块中的编号例如启动add_kernel100,256();代表 100 个 block每个 block 256 个线程。总线程100 × 256 25600。线程编号block0 thread0 i0 block0 thread1 i1 ... block1 thread0 i256这样可以处理大数组。完整版本__global__voidadd_kernel(float*a,float*b,float*c,intn){intiblockIdx.x*blockDim.xthreadIdx.x;if(in){c[i]a[i]b[i];}}解释计算线程编号intiblockIdx.x*blockDim.xthreadIdx.x;例如blockIdx.x 3 blockDim.x 256 threadIdx.x 10那么i 3 × 256 10 i 778这个线程负责a[778]b[778]边界检查if(in)防止线程数量 数据数量例如数据 1000 个启动 1024 个线程最后 24 个线程不能访问a[1000] a[1001] ...否则会越界错误。例子调用add_kernel2,4(a,b,c,n);表示 2 个 block每个 block 4 个线程blockIdx.xthreadIdx.x计算 i000011022033104115126137所以 8 个线程负责c[0] c[1] ... c[7]其中 CUDA Runtime 在创建线程时会自动给每个线程分配内置编号threadIdx.x。也就是说线程编号不是你从100,256算出来的而是 CUDA 运行时创建线程时自动维护的。
详解 CUDA 算子
1. 什么是算子一个算子Operator / Op就是一个基本计算节点。例如PyTorch 代码c torch.matmul(a, b)从用户角度看是两个矩阵相乘但是 PyTorch 内部会变成MatMul Operator → CUDA Kernel → GPU 执行这里MatMul 是一个算子CUDA Kernel 是真正跑在 GPU 上的代码2. CUDA 算子和 CUDA Kernel 的关系Kernel 是 CUDA 中最底层的 GPU 函数。例如__global__voidadd_kernel(float*a,float*b,float*c){intithreadIdx.x;c[i]a[i]b[i];}调用方式add_kernel...();上述代码详解1.__global__是什么__global__voidadd_kernel()是 CUDA 特有的关键字表示这个函数是一个 CUDA Kernel可以由 CPU 调用但是运行在 GPU 上。2. threadIdx.xintithreadIdx.x;这是 CUDA 最重要的概念之一。GPU 不是一个 CPU 核心执行 for 循环而是很多线程同时运行。例如启动add_kernel1,5();意思是1 个 block5 个线程。GPU 创建Thread 0 Thread 1 Thread 2 Thread 3 Thread 4每个线程都有自己的threadIdx.x于是线程threadIdx.x线程 00线程 11线程 22线程 33线程 44所以线程 0i0线程 1i1线程 2i23. 核心计算c[i]a[i]b[i];由于每个线程有不同的 i线程 0c[0]a[0]b[0]计算11011线程 1c[1]a[1]b[1]计算22022线程 2c[2]a[2]b[2]计算33033最终GPU 线程 Thread 0 c[0]a[0]b[0] Thread 1 c[1]a[1]b[1] Thread 2 c[2]a[2]b[2] Thread3: c[3]a[3]b[3] Thread4: c[4]a[4]b[4]多线程工作原理当前代码intithreadIdx.x;只能处理一个 block 里的线程。比如add_kernel1,256();最多 256 个元素如果数组 1000000 个元素怎么办需要block grid的 CUDA 三层结构Grid | | --- Block 0 | | | Thread0 | Thread1 | --- Block 1 | | | Thread0 | Thread1 | --- Block 2所以真实的 CUDA 代码通常是intiblockIdx.x*blockDim.xthreadIdx.x;也就是计算全局线程编号线程编号 block编号 × 每个block线程数量 线程在线程块中的编号例如启动add_kernel100,256();代表 100 个 block每个 block 256 个线程。总线程100 × 256 25600。线程编号block0 thread0 i0 block0 thread1 i1 ... block1 thread0 i256这样可以处理大数组。完整版本__global__voidadd_kernel(float*a,float*b,float*c,intn){intiblockIdx.x*blockDim.xthreadIdx.x;if(in){c[i]a[i]b[i];}}解释计算线程编号intiblockIdx.x*blockDim.xthreadIdx.x;例如blockIdx.x 3 blockDim.x 256 threadIdx.x 10那么i 3 × 256 10 i 778这个线程负责a[778]b[778]边界检查if(in)防止线程数量 数据数量例如数据 1000 个启动 1024 个线程最后 24 个线程不能访问a[1000] a[1001] ...否则会越界错误。例子调用add_kernel2,4(a,b,c,n);表示 2 个 block每个 block 4 个线程blockIdx.xthreadIdx.x计算 i000011022033104115126137所以 8 个线程负责c[0] c[1] ... c[7]其中 CUDA Runtime 在创建线程时会自动给每个线程分配内置编号threadIdx.x。也就是说线程编号不是你从100,256算出来的而是 CUDA 运行时创建线程时自动维护的。
