• 描述：有部分ptx代码无法被hipcc编译器编译，网上也无法找到关于ptx如何移植给hipcc的教程或编程手册。我们要移植的代码里包含如

__device__ __forceinline__ void CopyFromGlobalToShared(half __restrict__ (*SharedPTR)[...],
                                                            ^
unsigned smem_int_ptr = __cvta_generic_to_shared(smem_ptr);
                        ^
"l"(global_ptr),
^

这样的，无法被hipcc编译的代码。前两个用于优化的宏定义可能可以被移除，但第三个在ptx代码中的“l”不知道怎么处理。

• 可以通过下面的步骤来复现：

//t1.cu
#include <iostream>
#include <hip/hip_runtime.h>
__global__ void test_kernel(int *a, int *b) {
  __shared__ int smem[32];
  unsigned long long as = (unsigned long long)(smem+threadIdx.x);
  as &= 0x0FFFFFFULL;
  asm volatile("cp.async.ca.shared.global [%0], [%1], 4;\n" :: "l"(as), "l"(a+threadIdx.x));
  asm volatile("cp.async.wait_all;\n" ::);
  b[threadIdx.x] = smem[threadIdx.x];
}

int main(){
  int *a, *b;
  hipMallocManaged(&a, 32*sizeof(int));
  hipMallocManaged(&b, 32*sizeof(int));
  for (int i = 0; i < 32; i++) {a[i] = i; b[i] = 0;}
  test_kernel<<<1,32>>>(a, b);
  hipDeviceSynchronize();
  for (int i = 0; i < 32; i++) std::cout << b[i] << ",";
  std::cout << std::endl;
}

• 编译命令为

hipcc -o t1 t1.cu -lineinfo

• 编译失败。结果为

(/public/home/ykeye/yezhr7/env2) [ykeye@login04 fp6_test]$ hipcc -o t1 t1.cu -arch=sm_80 -lineinfo
clang-14: warning: argument unused during compilation: '-arch=sm_80' [-Wunused-command-line-argument]
t1.cu:8:64: error: invalid input constraint 'l' in asm
  asm volatile("cp.async.ca.shared.global [%0], [%1], 4;\n" :: "l"(as), "l"(a+threadIdx.x));
                                                               ^
1 error generated when compiling for gfx906.

• 请问有没有编程手册可以参考？或者仅对这个报错，有没有什么建议？