Jak używać clanga do kompilowania kodu OpenCL do ptx?

Question

Clang 3.0 jest w stanie skompilować OpenCL do ptx i użyć narzędzia Nvidii do uruchomienia kodu ptx na GPU. Jak mogę to zrobić? Proszę, bądź konkretny.

Answer 1

Zobacz przykład Justin Holewinski's blog dla konkretnego przykładu lub this thread, aby uzyskać bardziej szczegółowe instrukcje i łącza do przykładów.

Answer 2

Oto krótki przewodnik, jak to zrobić z trunk trunk Clang (w tym momencie 3.4) i libclc. Zakładam, że posiadasz podstawową wiedzę na temat konfigurowania i kompilowania LLVM i Clang, dlatego właśnie podałem listę flag konfiguracji, które użyłem.

square.cl:

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

1. kompilacji llvm i dzyń z nvptx wsparcia:

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

2. Get libclc (git clone http://llvm.org/git/libclc.git) i skompilować.

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

Jeśli masz problem kompilacji to może trzeba rozwiązać kilka nagłówków w ./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. jądrze skompilować assember LLVM IR :

   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. Jądro łącza z wbudowane implementacje z libclc

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

3. kompilacji całkowicie połączonych llvm IR PTX

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

square.nvptx.s:

// 
    // 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

Przy obecnej wersji z LLVM (3.4), libclc i nvptx back-end, proces kompilacji nieco się zmienił.

Musisz jawnie powiedzieć serwerowi nvptx, którego interfejsu sterownika użyć; Twoje opcje to nvptx-nvidia-cuda lub nvptx-nvidia-nvcl (dla OpenCL) i ich 64-bitowe odpowiedniki nvptx64-nvidia-cuda lub nvptx64-nvidia-nvcl.

Wygenerowany kod .ptx różni się nieznacznie w zależności od wybranego interfejsu. W kodzie zestawu utworzonym dla interfejsu API sterownika CUDA, elementy wewnętrzne .global i .ptr są usuwane z funkcji wejścia, ale są wymagane przez OpenCL.I zostały zmodyfikowane kompilacji kroki Mikael nieznacznie do produkcji kodu, który można uruchomić z OpenCL hosta:

1. skompilować 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. jądro Link:

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

3. Kompilacja z Ptx:

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