CUDA优化技巧

本文最后更新于：3 天前

向量化访存指令在SMEM中的机制

When you store (or load) more than 4 bytes per thread, which is like saying more than 128 bytes per warp, the GPU does not issue a single transaction. The largest transaction size is 128 bytes.

上面这段话是在Nvidia论坛上官方回复的一句话，可以这么理解：所谓的访存请求是会被拆分为memory transaction来执行的，如果是每个线程取/存4B的元素，则整个warp是处理128B的数据，这正好是单次memory transaction可以处理的最大数据

如果单次访存请求超过了128B，则不会只发射1条transaction

GPU会进行拆分，比如你采用LDS.64访存，会被拆成两条memory transaction，[T0-T15],[T16,T31]各自承担1个memory transaction。每一个都是128B的位宽，而后的bank conflict是根据每个memory transaction中涉及到的线程所产生的，而不是所谓的请求/warp/指令。

LDS.32

这是最朴素的一种访存机制，因为在warp内的32个线程，每个线程去访问共享内存中的32bit，即4B的一个元素。存在以下情况：

  • 线程访问到各自对应的共享内存上的元素，即访问到不同bank，无bank conflict;
  • 某些个线程访问到同一个bank上的同一地址，发起广播机制（这某些个线程里：其中1个线程取数，然后广播给另几个线程），无bank conflict;
  • 某些个线程访问到同一个bank上的不同地址，产生bank conflict。所谓的某些个线程有n个，则为n way bank conflict，会将访存转换为串行的n条指令进行（应该就是n条memory transaction）

LDS.64

采用类似于float2/uint2的向量化访存对shared memory进行访存时，1个warp被拆成了2个half warp，如T0-T15、T16-T31，1个活跃的half warp进行1次memory transaction，活跃即意味着里面有至少一个线程是有访存需求的。

因此如果1个warp里的2个half warp都是活跃的，则要进行2次memory transaction，是串行执行的。

当有以下情况发生时，则2次memory transaction可以合并成1次：

  • 对于整个warp，活跃的第i号线程，与不活跃的第i^1号线程 || 活跃的第i号线程和活跃的第i^1号线程访问同一块地址;
  • 对于整个warp，活跃的第i号线程，与不活跃的第i^2号线程 || 活跃的第i号线程和活跃的第i^2号线程访问同一块地址;

注：上面的合并memory transaction的情况是针对整个warp而言的，如果两个half warp一个满足规则1（就是第一条），一个满足规则2，也是无法合并的

LDS.128

采用类似于float4/uint4的向量化访存对shared memory进行访问，1个half warp被拆分成2个quarter warp，如T0-T7、T8-T15，1个活跃的quarter warp进行1次memory transaction

可见当4个quarter warp都活跃时需要进行4次串行的memory transaction，它的合并策略和LDS.64的一样，这里重复一次，并注明它的注意事项：

  • 对整个warp：(活跃的第i号线程与不活跃的第i^1号线程) || (活跃的第i号线程和活跃的第i^1号线程均访问同一地址);
  • 对整个warp：(活跃的第i号线程与不活跃的第i^2号线程) || (活跃的第i号线程和活跃的第i^2号线程均访问同一地址);

注：

  • 只有T0-T7,T8-T15活跃，并且满足条件1，则合并为1个half warp，1次memory transaction
  • 只有T8-T15,T16-T23活跃，且满足条件1，由于分属不同half warp，无法合并，2次memory transaction
  • T0-T7,T8-15满足条件1，T16-T23,T24-T31满足条件2，无法合并，4次memory transaction
  • T0-T7,T8-T15,T16-T23,T24-T31均满足同一合并条件，则合并为两个half warp，2次memory transaction，注意，俩half warp不可再合了，没有这种cascade操作

实验部分

LDS.64

源码如下：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
#include <cuda_runtime.h>
#include <cstdio>
#include <cstdint>

__global__ void smem_t1(uint32_t* src){
    __shared__ uint32_t smem[128];
    int tid = threadIdx.x;
    for(int i=0;i<4;i++){
        smem[i*32+tid] = tid;
    }
    __syncwarp();
    if(tid<16){
        reinterpret_cast<uint2*>(src)[tid]= reinterpret_cast<uint2*>(smem)[tid];
    }
}

__global__ void smem_t2(uint32_t* src){
    __shared__ uint32_t smem[128];
    int tid = threadIdx.x;
    for(int i=0;i<4;i++){
        smem[i*32+tid] = tid;
    }
    __syncwarp();
    if(tid<15 || tid==16){
        reinterpret_cast<uint2*>(src)[tid] = reinterpret_cast<uint2*>(smem)[tid==16?15:tid];
    }
}

__global__ void smem_t3(uint32_t* src){
    __shared__ uint32_t smem[128];
    int tid = threadIdx.x;
    for(int i=0;i<4;i++){
        smem[i*32+tid] = tid;
    }
    __syncwarp();
    reinterpret_cast<uint2*>(src)[tid] = reinterpret_cast<uint2*>(smem)[tid/2];
}

__global__ void smem_t4(uint32_t *src){
    __shared__ uint32_t smem[128];
    uint32_t tid = threadIdx.x;
    for(int i=0;i<4;i++){
        smem[i*32+tid] = tid;
    }
    __syncwarp();
    uint32_t idx = (tid>>4)<<3;
    // printf("0:[tid:%d]:[idx:%d]\n",tid,idx);
    idx = (idx == 0) ? idx+tid>>1 : (idx + (tid&1) + (((tid>>2)-4)<<1));
    // printf("1:[tid:%d]:[idx:%d]\n",tid,idx);
    reinterpret_cast<uint2*>(src)[tid] = reinterpret_cast<uint2*>(smem)[idx];
}

__global__ void smem_t5(uint32_t* src){
    __shared__ uint32_t smem[128];
    int tid = threadIdx.x;
    for(int i=0;i<4;i++){
        smem[i*32 + tid] = tid;
    }
    __syncwarp();
    uint32_t idx = tid>>4==0 ? tid : tid-16;
    reinterpret_cast<uint2*>(src)[tid] = reinterpret_cast<uint2*>(smem)[idx];
}


int main(){
    uint32_t* dA;
    size_t sizedA = sizeof(uint32_t) * 128;
    /*
        相当于:
        bank    0   1  ... 31 
               | | | | | | | |
                -   -   -   - 
               | | | | | | | |
                -   -   -   -
               | | | | | | | |
                -   -   -   - 
               | | | | | | | |
     */
    cudaMalloc(&dA,sizedA);
    dim3 grid(1);
    dim3 block(32);

    smem_t1<<<grid,block>>>(dA);
    smem_t2<<<grid,block>>>(dA);
    smem_t3<<<grid,block>>>(dA);
    smem_t4<<<grid,block>>>(dA);
    smem_t5<<<grid,block>>>(dA);
    cudaDeviceSynchronize();

    cudaFree(dA);
    return 0;
}

这里是对着第三个参考文件里的图LDS.64的图写的，并采用Nsight Compute的shard_op_ld进行bank conflict和memory transaction次数的查看，实验结果如下(其中wavefront代表memory transaction的次数，这里没有产生bank conflict的例子)：

lds64_smem_t1 lds64_smem_t2 lds64_smem_t3 lds64_smem_t4 lds64_smem_t5

参考文件：

Ampere架构的异步拷贝

异步拷贝的作用是global mem -> shared mem而无需借助register中转，在说异步拷贝前，先需要说下state space这个玩意

关于state space的一些理解

在ptx/sass级别，cuda内部有一个状态空间(state space)的系统,这玩意是一个分区的寻址结构(给我一种ptx上的namespace的感觉),按照参考文件2的解释:

A state space is a storage area with particular characteristics.

状态空间这个存储区域有着自己的特性,所有的变量都是驻留在一些状态空间中的.状态空间的特性有:大小,可否寻址,访问速度,访问权限,以及线程间的共享级别

下图是所有的状态空间:

state-space

下图是各状态空间对应的特性:

properties of state space

说完state space,回到异步拷贝上来

我们在cuda中对其中一个state space进行访问可以使用以下两种类型的指针的任意一种：

  • 被修饰的适合于该空间的指针；
  • generic pointer，这种类型的指针没有携带任何的state或者额外的信息去声明它所属于的state space；

在CUDA C中，指针就是指针，没有额外的修饰或者说元信息,当使用c中公开的方法去进行异步拷贝，不需要显式的转换；

而当你使用PTX指令(比如内联汇编)去做比如异步拷贝的时候就需要一个shared state space pointer，而不是generic pointer。就比如你用c++取得这个需要的state space地址，然后用ptx内联汇编去发射这条指令.你可以用__cvta_generic_to_shared()来进行地址的转换,从而获得shared state space的pointer

在ptx/sass中的state space pointer在c++中没有可以用来模拟的东东，因此使用了size_t的返回类型来识别64位的数量,虽然好像共享内存state space的pointer它是32位的(?)

shared state space pointer seems to easily fit in 32 bits

cp.async指令详情

1
2
3
4
5
6
7
8
9
10
11
12
cp.async.ca.shared{::cta}.global{.level::cache_hint}{.level::prefetch_size}
                         [dst], [src], cp-size{, src-size}{, cache-policy} ;
cp.async.cg.shared{::cta}.global{.level::cache_hint}{.level::prefetch_size}
                         [dst], [src], 16{, src-size}{, cache-policy} ;
cp.async.ca.shared{::cta}.global{.level::cache_hint}{.level::prefetch_size}
                         [dst], [src], cp-size{, ignore-src}{, cache-policy} ;
cp.async.cg.shared{::cta}.global{.level::cache_hint}{.level::prefetch_size}
                         [dst], [src], 16{, ignore-src}{, cache-policy} ;

.level::cache_hint =     { .L2::cache_hint }
.level::prefetch_size =  { .L2::64B, .L2::128B, .L2::256B }
cp-size =                { 4, 8, 16 }
  • dstshared state space的地址
  • src: global state space的地址
  • cp-size: byte为单位的整型常量,表示的是拷贝的数据的大小,可选值为4,8,16

注: cg表示只开了L2Cache;ca表示还开了L1cache.这里的可选项.level::prefetch_size看起来是预取到L2Cache,有不同粒度的选择;另一个可选项没搞懂,看起来跟缓存策略有关(看官网解释开缓存策略那个.L2::cache_hint必须用上);而src-size应该是为了解决不满足cp-size情况下时使用,比如3字节,cp-size为4,则第4个字节填充0

因为是异步拷贝,异步是nonblock的,所以你要拿个东西阻塞它,如果需要在某些代码执行前获取到结果的话:

1
cp.async.wait_all;

也可以:

1
2
cp.async.commit_group;	// 标记同步点
cp.async.wait_group 0;	// 同步到特定的同步点.这个0实际上是一个整型常量的operand N,表示至多允许N个未完成的异步事务,此处为0,即所有异步事务都得完成

上下这俩wait的语句块是等价的

小实验

这里的小实验的目的是用异步拷贝,把A的值给它异步拷贝到共享内存去,然后再从共享内存放到B中去,其中:

  • A原有的值A: [0,1,2,3,4,5,...,31]
  • B期望的结果B:[31,30,29,28,...,0]

我们每个线程(这里就是只开了1个block且只有32个线程)负责1个int类型的数据的移动,恰好只有32个数,1个线程也只需要搬运1次~

1
2
3
4
5
6
7
8
9
__global__ void test_cp_async(int* __restrict__ A,int* __restrict__ B){
    int tid = threadIdx.x;
    __shared__ int smem[32];
    size_t smemAddr = __cvta_generic_to_shared(smem + tid);
    asm volatile("cp.async.ca.shared.global [%0],[%1],4; \n" ::"l"(smemAddr),"l"(A+31-tid));
    asm volatile("cp.async.commit_group; \n" ::);
    asm volatile("cp.async.wait_group 0;  \n" ::);
    B[tid] = smem[tid];
}

实验结果如下图所示:

test_cp_async

怎么就支持了GMEM->SMEM呢

参考文件：

smem bank conflict free的一些方式

写在这里了tensor-core学习

参考文件:

CUDA OpDev

本博客所有文章除特别声明外，均采用 CC BY-SA 4.0 协议 ，转载请注明出处！