使用 ptx 指令“ldmatrix”和“mma”编译 cuda 时出错

Question

当我打算使用ldmatrix 和mma 指令时，我从以下代码中得到了错误。 PTX Docu 说 PTX 6.5 中引入了“ldmatrix”。所以我怀疑 PTX 版本可能是原因之一。我想知道我们如何找出我们使用的是哪个 PTX 版本？这些错误的原因还有哪些其他可能性？

    __device__ void
    runldmatrix(typet & D, unsigned addr){
      #if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 750))
        int x, y, z, w;            
        asm volatile (
          "ldmatrix.sync.aligned.x4.m8n8.shared.b16 {%0, %1, %2, %3}, [%4];" 
          : "=r"(x), "=r"(y), "=r"(z), "=r"(w) 
          : "r"(addr));             
        reinterpret_cast<int4 &>(D) = make_int4(x, y, z, w);
      #else
        assert(0);
      #endif
    }

    __device__  void
    runmma(typet & d, typet const & a, 
                typet  const & b,typet const & c ){
                  
#if (defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 750))                  
        unsigned const *A = reinterpret_cast<unsigned const *>(&a);
        unsigned const & B = reinterpret_cast<unsigned const &>(b);
        unsigned const *C = reinterpret_cast<unsigned const *>(&c);
        unsigned *D = reinterpret_cast<unsigned *>(&d);
        
          asm volatile(
              "mma.sync.aligned.m16n8k8.row.col.f16.f16.f16.f16 {%0,%1}, {%2,%3}, {%4}, {%5,%6};\n"
              : "=r"(D[0]), "=r"(D[1])
              : "r"(A[0]), "r"(A[1]), "r"(B), "r"(C[0]), "r"(C[1]));
#endif    
    }

/tmp/tmpxft_00002eb6_00000000-5_test_gemm.ptx，第 2637 行；错误 ： 未知修饰符 '.x4' ptxas

/tmp/tmpxft_00002eb6_00000000-5_test_gemm.ptx，第 2637 行；错误 ： 未知修饰符 '.m8n8' ptxas

/tmp/tmpxft_00002eb6_00000000-5_test_gemm.ptx，第 2637 行；错误 ： 不是任何已知指令的名称：'ldmatrix'

ptxas /tmp/tmpxft_000

02ee1_00000000-5_test_gemm.ptx，第 2611 行；错误：未知修饰符 '.m16n8k8' ptxas /tmp/tmpxft_00002ee1_00000000-5_test_gemm.ptx，线 2611;错误：指令“mma”需要形状修饰符

更新：

我用的是2080 Ti和CUDA 10.1，下面的cmake保证7.5的计算能力

cmake_minimum_required(VERSION 3.18)
project(Hello)
enable_language(CUDA)

add_executable(gunne test_gemm.cu)
target_include_directories(gunne PRIVATE include)
set_property(TARGET gunne PROPERTY CUDA_ARCHITECTURES 75)

Answer 1

Docu 说 PTX 6.5 中引入了“ldmatrix”。所以我怀疑 PTX 版本可能是原因之一。

还有哪些可能是导致这些错误的原因？

其实就是这个原因。 CUDA 10.1（它的最新版本）包括 PTX 版本 6.4。

如果搜索the PTX manual that shipped with that version of CUDA，则没有ldmatrix 指令。

此外，如果我们查看 relevant section，我们会发现该 PTX 版本中没有 mma.sync.aligned 指令的 m16n8k8 变体。

您的所有编译错误似乎都归结为这些问题。当我为typet（和#include <cassert>）提供定义时，代码将为我编译，例如CUDA 11.4

我想知道我们如何知道我们使用的是哪个 PTX 版本？

我至少能想到几种方法，可能还有其他方法。

“离线”方法是：假设您使用的是 8.0 或更高版本的 CUDA，请转到 cuda docs page，选择 the PTX manual，然后注意顶部的符号：

PTX ISA (PDF) - v11.5.1（旧版）

单击older 链接，您将进入一个页面，您可以在其中选择与您的CUDA 版本相对应的版本化在线文档。然后选择那里的 PTX 手册，它会显示它的版本。

另一种方法是使用您的工具链将任何 CUDA 代码编译为 PTX（例如 nvcc my_kernel_code.cu --ptx 并研究生成的 ptx 文件。在顶部附近，它将具有如下符号：

.version 7.4

这将告诉您您的工具链正在生成的 PTX 版本。

我并不是说您的代码在其他方面是正确的，只是在使用合适的工具链时它可以/将会编译。您没有提供完整的代码，也没有说明您的意图，所以我认为超越这一点没有真正意义，但是 int 变量与 .b16 指令的使用对我来说没有多大意义。但是，它似乎可以编译。