我有一个OpenMP代码段(并行区域(。我想在指示的行添加内联PTX组件(如下所示(我使用的是Nvidia gtx1070gpu和ubuntu Linux以及clang编译器。我尝试了asm volatile()方法，但clang编译器抛出编译错误(如下所示(
例如，如果我在CUDA内核内调用以下函数warp_id((，我们可以提取当前线程执行的warp_id：
__forceinline__ __device__ unsigned warp_id()
{
unsigned ret;
asm volatile ("mov.u32 %0, %warpid;" : "=r"(ret));
return ret;
}
我想做同样的事情，但在一个OpenMP并行区域内。
#pragma omp target teams distribute num_teams(15)
for(int i=0; i<R; i++)
{
#pragma omp parallel for default(shared) schedule(auto)
for(int j = 0; j < C; j++)
{
// extract warpid of this current thread iteration
}
}
Clang错误：
"/usr/local/cuda-9.2/bin/ptxas" -m64 -O3 -v --gpu-name sm_61 --output-file /tmp/openmp_ptxasm-1fb0b2.cubin /tmp/openmp_ptxasm-07f1f7.s -c
ptxas info : 59 bytes gmem
ptxas info : Compiling entry function '__omp_offloading_817_6600b8_main_l32' for 'sm_61'
ptxas info : Function properties for __omp_offloading_817_6600b8_main_l32
56 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info : Used 72 registers, 416 bytes cmem[0], 4 bytes cmem[2]
ptxas info : Function properties for _ZN27omptarget_nvptx_LoopSupportIllE13dispatch_nextEPiPlS2_S2_
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
"/usr/local/cuda-9.2/bin/nvlink" -o /tmp/openmp_ptxasm-0376ca.out -v -arch sm_61 -L. -L/home/mglab1/Desktop/llvm_new/install/lib -L/usr/local/cuda-9.2/lib64 -L/home/mglab1/Desktop/llvm_new/build/lib -lomptarget-nvptx /tmp/openmp_ptxasm-1fb0b2.cubin
nvlink info : 703720068 bytes gmem
nvlink info : Function properties for '__omp_offloading_817_6600b8_main_l32':
nvlink info : used 72 registers, 56 stack, 1084 bytes smem, 416 bytes cmem[0], 4 bytes cmem[2], 0 bytes lmem
"/home/mglab1/Desktop/llvm_new/build/bin/clang-7" -cc1 -triple x86_64-unknown-linux-gnu -emit-obj -disable-free -disable-llvm-verifier -discard-value-names -main-file-name openmp_ptxasm.cpp -mrelocation-model static -mthread-model posix -fmath-errno -masm-verbose -mconstructor-aliases -munwind-tables -fuse-init-array -target-cpu x86-64 -dwarf-column-info -debugger-tuning=gdb -momit-leaf-frame-pointer -v -resource-dir /home/mglab1/Desktop/llvm_new/build/lib/clang/7.0.0 -O3 -Wall -fdebug-compilation-dir /home/mglab1/Codes/cuda -ferror-limit 19 -fmessage-length 80 -fopenmp -fobjc-runtime=gcc -fdiagnostics-show-option -fcolor-diagnostics -vectorize-loops -vectorize-slp -o /tmp/openmp_ptxasm-34c034.o -x ir /tmp/openmp_ptxasm-8ccc6a.bc -fopenmp-targets=nvptx64-nvidia-cuda -faddrsig
clang -cc1 version 7.0.0 based upon LLVM 7.0.0 default target x86_64-unknown-linux-gnu
<inline asm>:1:10: error: invalid register name
mov.u32 %0, %smid;
^~
error: cannot compile inline asm
1 error generated.
如果我在内部循环中使用内联PTX，就会产生这个错误，如下所示：
#pragma omp parallel for default(shared) schedule(auto)
for(int j = 0; j <C; j++)
{
int thread_id = omp_get_thread_num();
unsigned int wid;
asm volatile("mov.u32 %0, %%warpid;" : "=r"(wid));
printf("Iteration= c[ %d ][ %d ], Team=%d, Thread=%d, warp=%dn",n, j, team_id, thread_id, wid);
S[n][j] = A[n][j] * B[n][j];
}
编译cmd:crang++-fopenmp-fopenmp targets=nvptx64 nvidia-cuda-Xopenmp target-march=sm_61-墙-O3 openmp.cpp
由于clang编译此类代码的方式，以及clang和OpenMP都没有任何语法功能，可以让编译器知道asm应该被视为设备代码，我认为这是不可行的。使用OpenACC可能有这样的空间，但不使用OpenMP编译器驱动的加速器卸载。
[此答案由添加到另一个问题的评论和信息组成，OP删除了该问题，并将其添加为社区wiki条目，以将其从CUDA标签的未回答队列中删除。请随时编辑并投票支持]