【问题标题】:Question on tensor core GEMM implementation?关于张量核心 GEMM 实现的问题?
【发布时间】:2021-11-02 06:45:03
【问题描述】:

我正在阅读一些张量核心material 和简单 GEMM 上的相关代码。我有两个问题:

1、D=A*B+C使用张量核时,将两个fp16矩阵4x4相乘,并将乘积fp32矩阵加到fp32累加器中。为什么两个fp16输入乘法A*Bresults为fp32类型?

2、在代码example中,为什么需要比例因子alphabeta?在示例中,它们被设置为2.0f

来自 NV 博客的代码 sn-p:

for(int i=0; i < c_frag.num_elements; i++) {
         c_frag.x[i] = alpha * acc_frag.x[i] + beta * c_frag.x[i];
      }

【问题讨论】:

    标签: cuda


    【解决方案1】:
    1. 在这种情况下,Tensorcore 设计人员选择提供 FP32 累加选项,以便可以以更高的精度(更多的尾数位)和更大的范围(更多的指数位)来表示许多乘法累加步骤的结果。这被认为对于他们想要支持的整体计算问题很有价值,包括 HPC 和 AI 计算。两个 FP16 数字的乘积可能无法在 FP16 中表示,而更多或大多数两个 FP16 数字的乘积可以在 FP32 中表示。

    2. 提供了比例因子alphabeta,以便提供的GEMM操作可以轻松对应众所周知的BLASGEMM操作,该操作广泛用于数值计算。这使开发人员可以更轻松地使用 Tensorcore 功能在现有数值计算代码中提供常用的计算范式。 CUBLAS GEMM 实现提供这些可调参数也是同样的原因。

    【讨论】:

    • 嗨罗伯特,我有一个关于这个linewmma::mma_sync(acc_frag, a_frag, b_frag, acc_frag);的问题,为什么它需要4个参数?我认为acc_frag = a_frag*b_frag,在什么情况下我们需要另一个acc_frag?抱歉,我没有找到任何相关文件。
    • 还记得你说的函数实现了什么吗?在你的问题中,你说:“当使用张量核心 D=A*B+C”时,D、A、B、C 有 4 个参数。在这个特定示例中两个相同的原因是我们要求它计算C=A*B+C,这很好。见here。这是编程指南的一部分,它可能是学习者最重要的文档。您可以通过在该文档中搜索 mma_sync 来快速找到它。可能对未来的研究有用。
    猜你喜欢
    • 2012-04-15
    • 1970-01-01
    • 1970-01-01
    • 2021-06-11
    • 1970-01-01
    • 2021-12-06
    • 2011-03-20
    • 1970-01-01
    • 1970-01-01
    相关资源
    最近更新 更多