记录CUTLASS的low-bit矩阵乘法GEMM计算中,主机内存到设备内存的搬运过程。(简单版)
(防止自己忘记。)
一、float型
首先对于float型的输入矩阵A和B而言,在主机中使用cutlass gemm模版创建两个tenosr来装矩阵A和B,然后通过sync_device函数将主机中的矩阵A和B放进设备中。在这个过程中,主机中的矩阵A和B的内存大小等于在设备中的矩阵A和B的内存大小。最后进行矩阵A和B的矩阵乘法计算。
二、low-bit型
当输入矩阵A和B是low-bit数据类型时,会有一些小变化。
例如,当输入矩阵A和B是1bit类型时,且假设矩阵A和矩阵B大小分别是8*128(row-major),128*8(col-major)。这里是指假设1bit为一个数据,那么矩阵A有8*128个数据,矩阵B有128*8个数据。
如果使用gemm模版来计算矩阵乘法时,首先会使用
cutlass::HostTensor<typename Gemm::ElementA, typename Gemm::LayoutA> tensor_A; cutlass::HostTensor<typename Gemm::ElementB, typename Gemm::LayoutB> tensor_B;
在主机中生成矩阵A和矩阵B,此时矩阵A和B的数据个数应该各自是8*128,128*8。但是由于CUTLASS中装载1bit数据的数据类型是CUTLASS自己定义的cutlass::uint1b_t,这个数据类型实际上是由uint8改装而成的。因此此时在主机中,矩阵A和矩阵B各自的内存大小应该是8*128*8bit,128*8*8bit。也就是说,每一个uint8数据类型中,就只有1个有效数据。
注意,其中的ElementA和ElementB都是指的矩阵A和B的各自的数据类型,在这里应该是
cutlass::uint1b_t (这个数据类型的特点以后再记录)。
在主机中完成1bit矩阵A和B的定义和生成后,需要将他们使用sync_device函数传输到GPU设备中,进行实际的矩阵乘法计算。
在传输到GPU设备后,矩阵A和B的实际内存大小变成了8*128bit,128*8bit。
这说明,在进行sync_device这个命令时,将矩阵A和B从主机传输到设备前,会先将矩阵A和B压缩。具体压缩过程是,将主机中的每一个cutlass::uint1b_t(uint8)的一个有效bit位提取出来,然后每8个为一组,放进GPU设备的cutlass::uint1b_t(uint8)中。所以矩阵A在设备中,应该就是有128个cutlass::uint1b_t(uint8)数据类型的数据。矩阵B同理。
如果,1bit矩阵A和B已经在GPU设备中了,那么只需要将1bit矩阵A和B的有效数据压缩到cutlass::uint1b_t(uint8)中,然后将数据传递到gemm矩阵计算模版中,就可以计算1bit矩阵乘法了。
(如有错误,请不吝指正,万分感谢。)