Как вставить встроенный PTX в код OpenMP? Как мы можем сделать в коде CUDA для сборки Nvidia PTX? - PullRequest
Новая
       137

Как вставить встроенный PTX в код OpenMP? Как мы можем сделать в коде CUDA для сборки Nvidia PTX?

2 голосов
/ 01 мая 2020

У меня есть этот сегмент кода OpenMP. Я хочу добавить встроенную сборку PTX в указанной строке. Я использую NVIDIA GTX1070 GPU и Ubuntu Linux и Clang компилятор. Я пробовал метод asm volatile (), но компилятор clang выдает ошибку компиляции (показано ниже) 

 #pragma omp target data map(to: A,B) map(from: S)
{
    #pragma omp target teams distribute  parallel for num_teams(15) \
     default(shared) schedule(auto) collapse(2)
    for(int n=0; n<R; n++)
    {

    for(int j = 0; j <C; j++)
    {
        int thread_id = omp_get_thread_num();
        int team_id= omp_get_team_num();  
        limits = omp_get_num_teams();   

        // ADD ASM PTX to print which SM is running this thread

            printf("Iteration= c[ %d ][ %d ], Team=%d, Thread=%d, Warp_id=%d\n",n, j, team_id, thread_id,warp_id);
        S[n][j] = A[n][j] * B[n][j];
    }

    }
}

Ошибка 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 в внутренний l oop как это: 

#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=%d\n",n, j, team_id, thread_id, wid);
        S[n][j] = A[n][j] * B[n][j];
    }
...