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

Question

当我打算使用 ldmatrix 和 mma 指令时，我从以下代码中得到了错误。 PTX Docu 说 'ldmatrix' 是在 PTX 6.5 中引入的。所以我怀疑 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, line 2637; error : Unknown modifier '.x4' ptxas

/tmp/tmpxft_00002eb6_00000000-5_test_gemm.ptx, line 2637; error : Unknown modifier '.m8n8' ptxas

/tmp/tmpxft_00002eb6_00000000-5_test_gemm.ptx, line 2637; error : Not a name of any known instruction: 'ldmatrix'

ptxas /tmp/tmpxft_000

02ee1_00000000-5_test_gemm.ptx, line 2611; error : Unknown modifier '.m16n8k8' ptxas /tmp/tmpxft_00002ee1_00000000-5_test_gemm.ptx, line 2611; error : Shape modifier required for instruction '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 says that 'ldmatrix' is introduced in PTX 6.5. So I doubt the the PTX version could be one reason.

What other possibilities could be the reason for those errors?

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

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

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

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

I'd like to know how do we find out which PTX version are we using?

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

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

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

单击 older link，它将带您到一个页面，您可以在其中 select 与您的 CUDA 版本相对应的版本化在线文档。然后select那里的PTX手册，它会指出它是什么版本。

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

.version 7.4

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

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

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

Error when compile cuda with ptx instruction 'ldmatrix' and 'mma'

cmake

cuda

ptx