CUDA优化

CUDA优化

避免Warp分化

if (con)
{
    //do something
}
else
{
    //do something
}

假设这段代码是核函数的一部分，那么当一个线程束的32个线程执行这段代码的时候，如果其中16个执行if中的代码段，而另外16个执行else中的代码块，同一个线程束中的线程，执行不同的指令，这叫做线程束的分化。

解决bank conflict

没4个字节一个bank，一共有32个bank，然后每个bank有多个字段

在共享内存中，就像电影院的座位一样：一列的座位就相当于一个bank，所以每行有32个座位，在每个座位上可以“坐”一个32-bits的数据(或者多个小于32-bits的数据，如4个char型的数据，2个short型的数据)；而正常情况下，我们是按照先坐完一行再坐下一行的顺序来坐座位的，在shared memory中地址映射的方式也是这样的

其中0-31为bank编号,如果申请一个共享内存数组__shared__ int cache[64],int 恰好为4个字节，那么cache[0]访问bank[0][0], cache[1]访问bank[0][1],...,cache[31]访问bank[0][31],cache超过32时，cache就会去访问下一行的bank,即cache[32]就会访问bank[1][0],以此类推。

bank冲突就是在这样的条件下产生，即如果一个warp的多个线程访问同一个bank的不同字段时(注：不同字段如bank[0][0],bank[1][0],...,bank[n][0])，那么就发生了bank冲突，因为不同bank可以同时访问，而当如果多个线程请求的内存地址被映射到了同一个bank上，那么这些请求就变成了串行的。

利用空闲线程

展开最后一维减少同步

完全循环展开

OpenCore使用

调整block大小

shuffle指令

NV提出了Shuffle指令，对于reduce优化有着非常好的效果。目前绝大多数访存类算子，像是softmax，batch_norm，reduce等，都是用Shuffle实现。所以，在这里谈一下这么把shuffle指令用在reduce优化上。

Shuffle指令是一组针对warp的指令。Shuffle指令最重要的特性就是warp内的寄存器可以相互访问。在没有shuffle指令的时候，各个线程在进行通信时只能通过shared memory来访问彼此的寄存器。而采用了shuffle指令之后，warp内的线程可以直接对其他线程的寄存器进行访存。通过这种方式可以减少访存的延时。除此之外，带来的最大好处就是可编程性提高了，在某些场景下，就不用shared memory了。毕竟，开发者要自己去控制 shared memory还是挺麻烦的一个事。

关于shuffle指令见：shuffle

L1 cache

L1 cache 同样需要合并访存才能发挥最大带宽
L1 cache 是用于缓存全局内存和局部内存的缓存，其内存读取细粒度是128 bytes，所以内存对齐的访问地址最好是128 bytes的倍数；且访问的数据最好是一次128 bytes(100%带宽利用率)， 再大也只能逐次次访问，因为L1 cache line是128 bytes

nvcc -g -G -Xptxas -dlcm=ca -arch=sm_37 test.cu -o test

-Xptxas -dlcm=ca 开启L1 Cache
-Xptxas -dlcm=cg 禁用L1 Cache

向量化的访存

可以通过使用向量 load 指令来一次读取多个数据，float4 float2 int4 等数据类型的使用，nvcc编译器会在指令选择的时候选择向量 load 指令

__global__ void vectorizedMemoryAccess(float* input, float* output, int size) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int idx = tid * 4;  // float4向量化访存

    if (idx < size) {
        float4 data = reinterpret_cast<float4*>(input)[tid];
        float sum = data.x + data.y + data.z + data.w;
        output[tid] = sum;
    }
}