hip程序随手可加的小优化

hip程序随手可加的小优化

下面所说的优化，不会有质变的效果。但是，节省的一点点时间也会产生较大的加速比。而且很方便添加，在写kernel时可以随手加上。

向量化访存（正优化）

hip 提供了float2,float4用于向量化访存，如果我们想使用更长的向量化访存应该怎么做呢。

typedef float float8 __attribute__((ext_vector_type(8))); // 定义一个向量类型
#define FLOAT8(pointer) (reinterpret_cast<float8*>((void *)&(pointer))[0]) // 从此位置得到一个float8的变量
// 如何使用
float8 reg_8 = FLOAT8(d_A[0]); // d_A为全局内存指针，d_A的类型为float *
// 这样reg_8 就保存了8个 float 

其实底层仅仅提供了，float4的向量访存指令。float8会告诉编译器，我要访存8个float。编译器会更加方便去优化。float8比两个float4访存会快一些。

restrict 修饰符(正优化）

为所有的指针变量添加__restrict__ 修饰符，可以让核函数得到加速。__restrict__的详细介绍

__global__ void kernel(float* a, float* b, float* c) {
...
}

__global__ void kernel(float* __restrict__ a, float* __restrict__ b, float* __restrict__ c) {
...
}

launch_bounds 修饰符（第一位置的数字正优化，第二位置的数字正负优化都有可能）

该修饰符用于修饰核函数。为核函数添加一些限制，以便让编译器得到最优的代码。

使用方法一般如下：

__global__ __launch_bounds__(256, 2) void gemm_kernel(float* __restrict__ a, float* __restrict__ b, float* __restrict__ c) {

}

解释一下变量的含义：

“__ launch_bounds __(256, 2)” 中的 “256” 表示该内核中的一个block最多使用256个线程。如果我们不加这个限制，编译器就会按照默认最大的线程数（1024）来优化代码。这样，每个线程分到的寄存器会变少，影响性能。如果我们确定内核每个block最多使用256个线程，就可以加这个限制。性能会有所优化。

“__ launch_bounds __(256, 2)” 中的 “2” 表示一个CU中最少能同时运行几个block。如果加上了这个变量，编译器就会限制每个block所使用的资源。保证一个CU能同时运行两个block。这个改变对不同的内核有不同的效果。可以看情况添加。

__builtin_amdgcn_readfirstlane（正优化）

使用方法：

__global__ __launch_bounds__(256, 2) void gemm_kernel(float* __restrict__ a, float* __restrict__ b, float* __restrict__ c) {
      const index_t lds_ldb = n + padding; // 原来的方式
      const index_t lds_ldb = __builtin_amdgcn_readfirstlane(n + padding); // 添加 __builtin_amdgcn_readfirstlane
}

该指令表示同一个warp中的变量都是同一个值。其他的线程会从线程束的第一个线程中获取数据。

大家可能会想这样并没有节约时间，从第一个线程获取数据可能还会有通信开销。其实这个优化主要优化寄存器的个数。

在AMD中寄存器分为两种VGPR和SGPR。VGPR按照线程分配，SGPR按照warp分配。

通过__builtin_amdgcn_readfirstlane作用的变量会放到SGPR。放在VGPR中会使用64个寄存器（每个线程都会占用一个）。一般编译器会分析出哪些变量不会变，放置在SGPR中。有些变量我们自己知道一个warp中是不会变的，这时就可以通过__builtin_amdgcn_readfirstlane来放置到SGPR寄存器中。比如：warp的序号。示例如下：

const index_t wid = _builtin_amdgcn_readfirstlane(threadIdx.x / WARP_SIZE);

一般情况，编译器是分析不出这个变量在一个warp中是相同的，这时就需要我们手动添加。