CUTLASS的GEMM(low-bit 内存搬运过程记录)(简单版)

记录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矩阵乘法了。

(如有错误,请不吝指正,万分感谢。)

  • 6
    点赞
  • 5
    收藏
    觉得还不错? 一键收藏
  • 0
    评论

“相关推荐”对你有帮助么?

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值