我正在寻找提高基于OpenCl的算法效率的方法。
目前，我在Radeon VII卡上使用float和int数据类型。但是，包含-8到+7之间数字的数据类型就足够了。
根据下面的文章，Radeon VII在限制为INT8/INT4时实现了53/110 TFlops的峰值性能，这比float的14 TFlops高得多。
https://www.pcgameshardware.de/Radeon-VII-Grafikkarte-268194/Tests/Benchmark-Review-1274185/2/
所以我的问题是如何利用INT8/4运算？在OpenCl中简单地使用数据类型char而不是int？既然char是最小的内置数据类型，我怎么能使用INT4呢？
对于"int8"，即8位整数，OpenCL类型实际上是char(有符号，-128到+127(或uchar(无符号，0到255(不要与OpenCL类型int8混淆，后者是一个由8个32位整数组成的向量
为了获得良好的性能，您可能希望使用这些的矢量版本，如char4或char16，尽管这应该基于您的性能测量，而不是猜测。
请注意，您需要注意溢出行为，尤其是对于乘法，您可能需要对16位值执行中间运算。(short/ushort/short4/ushort16/等(OpenCL还提供了"饱和"加法和减法以及其他一些有用的整数内置函数。
我不知道OpenCL或任何其他GPGPU框架，甚至任何扩展中有任何对压缩4位整数数学的"原生"支持。也许有经验的人可以参与进来，但我的猜测是，你实际上需要使用位移和掩码来解包uchar值，对uchar值执行操作，然后打包回4位半字节进行存储。速度的提高可能来自于这样一个事实，即您可以使用8位逻辑安全地相乘，而不需要16位来捕获溢出。
我对一些内核进行了测试，看看int8和char8:之间的性能是否有差异
typedef int8 type_msg;
//typedef char8 type_msg;
#define convert_type_msg(x) convert_int8(x)
__kernel void some_operation(__global type_msg *in_buff,
__global type_msg *out_buff)
{
out_buff[get_global_id(0)] = in_buff[get_global_id(0)] +(type_msg)(2);
}
首先，为了了解GPU上发生了什么，我使用CodeXL来获得汇编代码。
以下是使用int8的汇编代码的一部分：
global_load_dwordx4 v[4:7], v[2:3], off
global_load_dwordx4 v[8:11], v[2:3], off inst_offset:16
v_add_co_u32 v0, vcc, s6, v0
v_mov_b32 v2, s7
v_addc_co_u32 v1, vcc, v2, v1, vcc
s_waitcnt vmcnt(0)
v_add_u32 v8, 2, v8
v_add_u32 v9, 2, v9
v_add_u32 v10, 2, v10
v_add_u32 v11, 2, v11
global_store_dwordx4 v[0:1], v[8:11], off inst_offset
v_add_u32 v2, 2, v4
v_add_u32 v3, 2, v5
v_add_u32 v4, 2, v6
v_add_u32 v5, 2, v7
global_store_dwordx4 v[0:1], v[2:5], off
这里是使用char8的汇编代码的一部分：
global_load_dwordx2 v[2:3], v[2:3], off
s_waitcnt vmcnt(0)
v_lshlrev_b32 v4, 8, v3 src1_sel:BYTE_3
v_lshrrev_b32 v5, 8, v3
v_add_u32 v6, 2, v3 src1_sel:WORD_1
v_add_u32 v4, 0x00000200, v4
s_movk_i32 s0, 0x00ff
v_lshlrev_b32 v7, 8, v2 src1_sel:BYTE_3
v_add_u32 v5, 2, v5
v_bfi_b32 v4, s0, v6, v4
s_mov_b32 s1, 0x02010004
v_lshrrev_b32 v6, 8, v2
v_add_u32 v8, 2, v2 src1_sel:WORD_1
v_add_u32 v7, 0x00000200, v7
v_add_u32 v3, 2, v3
v_perm_b32 v4, v5, v4, s1
v_add_u32 v5, 2, v6
v_bfi_b32 v6, s0, v8, v7
v_add_co_u32 v0, vcc, s6, v0
v_mov_b32 v7, s7
v_addc_co_u32 v1, vcc, v7, v1, vcc
v_perm_b32 v3, v3, v4, s1
v_add_u32 v2, 2, v2
v_perm_b32 v4, v5, v6, s1
v_perm_b32 v2, v2, v4, s1
global_store_dword v[0:1], v3, off inst_offset:4
global_store_dword v[0:1], v2, off
我不是汇编语言方面的专家，但据我所知，在这两种情况下，都使用v_add_u32操作进行了8次添加。此外，CCD_ 23似乎需要更多的操作，如CCD_ 24和CCD_ 25。也许有人可以解释一下他们在做什么。
使用char8的唯一好处似乎是，需要更少的全局内存访问。例如，对于char8只有一个global_load_dwordx2接入，但是对于int8有两个global_load_dwordx4接入。
因此，就性能而言，char8对于计算有界算法可能慢一点，但对于内存有界算法则快一点。
为了验证分析，我建立了一个小实验，其中算术是瓶颈。为了确保编译器不会过于简化for循环，我在其中添加了一些分支
typedef int8 type_msg;
#define convert_type_msg(x) convert_int8(x)
//typedef char8 type_msg;
//#define convert_type_msg(x) convert_char8(x)
__kernel void some_complex_operation(__global char8 *in_buff,
__global char8 *out_buff)
{
type_msg res = in_buff[get_global_id(0)];
for(int i=0; i<1000000; i++)
{
res += select((type_msg)(-1), (type_msg)(4), res<(type_msg)100);
}
out_buff[get_global_id(0)] =(type_msg) res;
}
在我的系统上的平均时间(运行100次(
int8为0.0558秒char8为0.0754秒short8为0.0738秒long8为0.1105秒
因此char8大约多花35%的时间。这证实了在汇编语言中为char8生成更多指令的观察结果。不过，对附加程序集语句进行一些专业的解释会很好。