用ptx指令'ldmatrix‘和'mma’编译cuda时出错
当我打算使用ldmatrix和mma指令时，我从下面的代码中得到了错误。PTX说'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，第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第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)回答 1
Stack Overflow用户
发布于 2022-01-01 15:03:13
Docu说，在PTX 6.5中引入了“ldmatrix”。所以我怀疑PTX版本可能是其中一个原因。造成这些错误的其他可能原因是什么？
事实上，这就是原因。CUDA 10.1 (它的最新版本)包括PTX版本6.4。
如果搜索该版本的CUDA附带的PTX手册，则没有ldmatrix指令。
此外，如果我们查看相关部分，我们会发现在PTX版本中，mma.sync.aligned指令中没有m16n8k8变体。
似乎所有的编译错误都归结到了这些问题上。当我为typet (和#include <cassert>)提供定义时，代码将在例如CUDA 11.4上为我编译
我想知道怎样才能找到我们使用的PTX版本？
我至少能想到几种方法，也许还有其他的方法。
“脱机”方法是:假设您使用的是CUDA版本的8.0或更高版本，转到库达文档页，选择PTX手册，然后在表示法的顶部注意：
PTX ISA (PDF) - v11.5.1 (旧版)
单击更老的链接，它将带您到一个页面，在该页面中，您可以选择与您的CUDA版本相对应的版本在线文档。然后在那里选择PTX手册，它将指明它的版本。
另一种方法是使用工具链将任何CUDA代码编译成PTX (例如，nvcc my_kernel_code.cu --ptx并研究生成的ptx文件)。在顶部附近将有一个符号，如：
.version 7.4这将告诉您您的工具链正在生成的PTX版本。
我并不是说您的代码在其他方面是正确的，只是说它在使用适当的工具链时可以/将编译。您没有提供完整的代码，也没有说明您的意图，因此我认为没有什么真正的意义要超越这一点，但是int变量与.b16指令的使用对我来说没有多大意义。然而，它似乎是在编译。
https://stackoverflow.com/questions/70549137
复制相似问题
腾讯云开发者
