我正在寻找提高基于OpenCl的算法效率的方法。
当前，我在Radeon VII卡上使用float和int数据类型。但是，覆盖-8至+7之间数字的数据类型就足够了。
根据以下文章，当限制为INT8 / INT4时，Radeon VII的峰值性能达到53/110 TFlops，比float的14 TFlops高得多。
https://www.pcgameshardware.de/Radeon-VII-Grafikkarte-268194/Tests/Benchmark-Review-1274185/2/amp/
所以我的问题是如何使用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还提供“饱和”加减法和a few other helpful integer built-in functions。
我不知道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次加法。因此，就性能而言，至少从组装者的角度来看，不应有差异。同样，char8似乎需要更多的操作，例如v_perm_b32和v_bfi_b32。也许有些人可以解释这些操作。
使用char8的唯一好处似乎是，需要较少的全局内存访问。例如。 global_load_dwordx2只有一个char8访问，而global_load_dwordx4只有2个int8访问。
为了验证分析，我建立了一个小实验，算术是瓶颈。为了使编译器不会简化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<100000; 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秒因此char8会消耗大约35％的时间。这证实了这一发现，即以汇编语言为char8生成了更多指令。但是，对附加的汇编语句进行一些专业的解释会很好。