clang을 사용하여 OpenCL을 ptx 코드로 컴파일하는 방법?

Question

Clang 3.0은 OpenCL을 ptx로 컴파일하고 Nvidia의 도구를 사용하여 GPU에서 ptx 코드를 실행할 수 있습니다. 어떻게해야합니까? 구체적으로 기재하십시오.

Answer 1

구체적인 예를 보려면 Justin Holewinski's blog을, 샘플에 대한 자세한 단계 및 링크는 this thread을 참조하십시오.

Answer 2

다음은 Clang 트렁크 (이 시점에서 3.4)와 libclc를 사용하여 수행하는 방법에 대한 간략한 안내서입니다. LLVM과 Clang을 구성하고 컴파일하는 방법에 대한 기본적인 지식이 있다고 가정하고, 필자가 사용했던 configure 플래그를 나열했습니다. nvptx 지원

__kernel void vector_square(__global float4* input, __global float4* output) { 
    int i = get_global_id(0); 
    output[i] = input[i]*input[i]; 
} 

1. 컴파일 LLVM과 그 소리 :

   square.cl

   ../llvm-trunk/configure --prefix=$PWD/../install-trunk --enable-debug-runtime --enable-jit --enable-targets=x86,x86_64,nvptx 
   make install 
   

2. (자식 클론 http://llvm.org/git/libclc.git)를 libclc 가져 오기 및 컴파일. 당신은 문제가이 컴파일이있는 경우

   ./configure.py --with-llvm-config=$PWD/../install-trunk/bin/llvm-config 
   make 
   

당신은 LLVM IR의 assember에 ./utils/prepare-builtins.cpp

-#include "llvm/Function.h" 
-#include "llvm/GlobalVariable.h" 
-#include "llvm/LLVMContext.h" 
-#include "llvm/Module.h" 
+#include "llvm/IR/Function.h" 
+#include "llvm/IR/GlobalVariable.h" 
+#include "llvm/IR/LLVMContext.h" 
+#include "llvm/IR/Module.h" 

1. 컴파일 커널 헤더의 몇 가지 문제를 해결해야 할 수도 있습니다 :

   clang -Dcl_clang_storage_class_specifiers -isystem libclc/generic/include -include clc/clc.h -target nvptx -xcl square.cl -emit-llvm -S -o square.ll 
   

2. llvm-link libclc/nvptx--nvidiacl/lib/builtins.bc square.ll -o square.linked.bc 
   

3. 컴파일이 완전히 PTX

   에 LLVM IR를 연결 libclc

   clang -target nvptx square.linked.bc -S -o square.nvptx.s 
   

square.nvptx.s에서 내장 구현 : LLVM의 현재 버전

// 
    // Generated by LLVM NVPTX Back-End 
    // 
    .version 3.1 
    .target sm_20, texmode_independent 
    .address_size 32 

      // .globl  vector_square 

    .entry vector_square(
      .param .u32 .ptr .global .align 16 vector_square_param_0, 
      .param .u32 .ptr .global .align 16 vector_square_param_1 
    ) 
    { 
      .reg .pred %p<396>; 
      .reg .s16 %rc<396>; 
      .reg .s16 %rs<396>; 
      .reg .s32 %r<396>; 
      .reg .s64 %rl<396>; 
      .reg .f32 %f<396>; 
      .reg .f64 %fl<396>; 

      ld.param.u32 %r0, [vector_square_param_0]; 
      mov.u32 %r1, %ctaid.x; 
      ld.param.u32 %r2, [vector_square_param_1]; 
      mov.u32 %r3, %ntid.x; 
      mov.u32 %r4, %tid.x; 
      mad.lo.s32  %r1, %r3, %r1, %r4; 
      shl.b32   %r1, %r1, 4; 
      add.s32   %r0, %r0, %r1; 
      ld.global.v4.f32  {%f0, %f1, %f2, %f3}, [%r0]; 
      mul.f32   %f0, %f0, %f0; 
      mul.f32   %f1, %f1, %f1; 
      mul.f32   %f2, %f2, %f2; 
      mul.f32   %f3, %f3, %f3; 
      add.s32   %r0, %r2, %r1; 
      st.global.f32 [%r0+12], %f3; 
      st.global.f32 [%r0+8], %f2; 
      st.global.f32 [%r0+4], %f1; 
      st.global.f32 [%r0], %f0; 
      ret; 
    } 

Answer 3

(3.4), libclc 및 nvptx 백엔드에서 컴파일 프로세스가 약간 변경되었습니다.

nvptx 백엔드에 사용할 드라이버 인터페이스를 명시 적으로 말해야합니다. nvptx-nvidia-cuda 또는 nvptx-nvidia-nvcl (OpenCL 용) 및 해당 64 비트에 해당하는 nvptx64-nvidia-cuda 또는 nvptx64-nvidia-nvcl이 있습니다.

생성 된 .ptx 코드는 선택한 인터페이스에 따라 약간 다릅니다. CUDA 드라이버 API 용으로 생성 된 어셈블리 코드에서 intrinsics .global 및 .ptr은 입력 함수에서 제외되지만 OpenCL에서는 필요합니다.

1. 컴파일 LLVM에 IR :

   clang -Dcl_clang_storage_class_specifiers -isystem libclc/generic/include -include clc/clc.h -target nvptx64-nvidia-nvcl -xcl test.cl -emit-llvm -S -o test.ll 
   

2. 링크 커널 :

   llvm-link libclc/built_libs/nvptx64--nvidiacl.bc test.ll -o test.linked.bc 
   

3. 나는 OpenCL을 호스트와 함께 실행할 수있는 코드를 생성하기 위해 약간 미카엘의 컴파일 단계를 수정 한 컴파일하여 Ptx :

   clang -target nvptx64-nvidia-nvcl test.linked.bc -S -o test.nvptx.s