1.算子开发主要就是cuda编程在这里为什么会使用cuda因为gpu相比于cpu有大量核心的小的处理器去处理相似或者相同的任务对于self attention的大量相似的运算非常合适。2.基础知识__global__ 标记是gpu上的代码然后在具体使用之间要 cudaMalloc、cudaMemcpy、调用、cudaMemcpy、cudaMalloc。关于grid和block的概念就是从启动kernel的理解程序是按照grid组织的一个grid包含多个block一个block包含多个线程同一个block中的线程是按照32个为单位去调度的每32个也叫做一个线程束即warp。在这里可以使用int作为grid和blocksize去传输也可以使用dim3去作为二维三维的去使用。3.矩阵乘法TILE加速主要是利用了共享内存提到存储模型有多种从线程独有的寄存器、到共享内存这个是片上的block共享还有全局内存这个是多个block都可以访问的更慢存储bound往往是因为这里。共享内存和cache的区别是共享内存是由程序员申请和管理一个block共享cache是硬件控制的有l1和l2l1是一个sm上的block共享而l2是多个sm上的线程共享。#define TILE__global__ func(float *A,float *B,float *C,int M,int K,int N){__shared__ float As[TILE][TILE];__shared__ float Bs[TILE][TILE];int rowblockDim.y*blockIdx.ythreadIdx.y;int colblockDim.x*blockIdx.xthreadIdx.x;float sum0.0f;for(int i0;i(KTILE-1)/TILE;i){int a_coli*TILEthreadIdx.x;int b_rowi*TILEthreadIdx.y;if(rowMa_colK){As[threadIdx.y][threadIdx.x]A[row*Ka_col];}else As[threadIdx.y][threadIdx.x]0;if(b_rowKcolN){Bs[threadIdx.y][threadIdx.x]B[b_row*Ncol];else Bs[threadIdx.y][threadIdx.x]0;__syncthreads();for(int k0;kTILE;k){sumAs[threadIdx.y][k]*Bs[k][threadIdx.x];}__syncthreads();}if(rowMcolN)C[row*Ncol]sum;}}调用方用二维去调用dim3 block(16,16);dim3 grid((col15)/16,(row15)/16);总结这里我的问题是二维矩阵不是指针本身就是指针、一次是计算TILE个不是K个。4.bank conflict的概念同一个共享内存分成32个bank在同一时刻不同线程不能访问同一个bank的不同地址这会导致线程延迟执行增加执行时间降低并行率解决方法可以在设计的时候避免stride是32的倍数、使用padding、使用寄存器直接交换数值的技术。
CUDA学习笔记总结
1.算子开发主要就是cuda编程在这里为什么会使用cuda因为gpu相比于cpu有大量核心的小的处理器去处理相似或者相同的任务对于self attention的大量相似的运算非常合适。2.基础知识__global__ 标记是gpu上的代码然后在具体使用之间要 cudaMalloc、cudaMemcpy、调用、cudaMemcpy、cudaMalloc。关于grid和block的概念就是从启动kernel的理解程序是按照grid组织的一个grid包含多个block一个block包含多个线程同一个block中的线程是按照32个为单位去调度的每32个也叫做一个线程束即warp。在这里可以使用int作为grid和blocksize去传输也可以使用dim3去作为二维三维的去使用。3.矩阵乘法TILE加速主要是利用了共享内存提到存储模型有多种从线程独有的寄存器、到共享内存这个是片上的block共享还有全局内存这个是多个block都可以访问的更慢存储bound往往是因为这里。共享内存和cache的区别是共享内存是由程序员申请和管理一个block共享cache是硬件控制的有l1和l2l1是一个sm上的block共享而l2是多个sm上的线程共享。#define TILE__global__ func(float *A,float *B,float *C,int M,int K,int N){__shared__ float As[TILE][TILE];__shared__ float Bs[TILE][TILE];int rowblockDim.y*blockIdx.ythreadIdx.y;int colblockDim.x*blockIdx.xthreadIdx.x;float sum0.0f;for(int i0;i(KTILE-1)/TILE;i){int a_coli*TILEthreadIdx.x;int b_rowi*TILEthreadIdx.y;if(rowMa_colK){As[threadIdx.y][threadIdx.x]A[row*Ka_col];}else As[threadIdx.y][threadIdx.x]0;if(b_rowKcolN){Bs[threadIdx.y][threadIdx.x]B[b_row*Ncol];else Bs[threadIdx.y][threadIdx.x]0;__syncthreads();for(int k0;kTILE;k){sumAs[threadIdx.y][k]*Bs[k][threadIdx.x];}__syncthreads();}if(rowMcolN)C[row*Ncol]sum;}}调用方用二维去调用dim3 block(16,16);dim3 grid((col15)/16,(row15)/16);总结这里我的问题是二维矩阵不是指针本身就是指针、一次是计算TILE个不是K个。4.bank conflict的概念同一个共享内存分成32个bank在同一时刻不同线程不能访问同一个bank的不同地址这会导致线程延迟执行增加执行时间降低并行率解决方法可以在设计的时候避免stride是32的倍数、使用padding、使用寄存器直接交换数值的技术。
