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和l2，l1是一个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 row=blockDim.y*blockIdx.y+threadIdx.y;
int col=blockDim.x*blockIdx.x+threadIdx.x;
float sum=0.0f;
for(int i=0;i<(K+TILE-1)/TILE;i++){
int a_col=i*TILE+threadIdx.x;
int b_row=i*TILE+threadIdx.y;
if(row<M&&a_col<K){
As[threadIdx.y][threadIdx.x]=A[row*K+a_col];
}else As[threadIdx.y][threadIdx.x]=0;
if(b_row<K&&col<N){
Bs[threadIdx.y][threadIdx.x]=B[b_row*N+col];
else Bs[threadIdx.y][threadIdx.x]=0;
__syncthreads();
for(int k=0;k<TILE;k++){
sum+=As[threadIdx.y][k]*Bs[k][threadIdx.x];
}
__syncthreads();
}
if(row<M&&col<N)C[row*N+col]=sum;
}
}
调用方用二维去调用
dim3 block(16,16);
dim3 grid((col+15)/16,(row+15)/16);
总结：这里我的问题是二维矩阵不是指针，本身就是指针、一次是计算TILE个，不是K个。
4.bank conflict的概念
同一个共享内存分成32个bank，在同一时刻不同线程不能访问同一个bank的不同地址，这会导致线程延迟执行增加执行时间降低并行率，解决方法可以在设计的时候避免stride是32的倍数、使用padding、使用寄存器直接交换数值的技术。
