Come usare clang per compilare OpenCL al codice ptx?

Question

Clang 3.0 è in grado di compilare OpenCL in ptx e utilizzare lo strumento di Nvidia per avviare il codice ptx su GPU. Come posso fare questo? Per favore sii specifico

Answer 1

Vedere Justin Holewinski's blog per un esempio specifico o this thread per alcuni passaggi più dettagliati e collegamenti agli esempi.

Answer 2

Ecco una breve guida su come farlo con il trunk Clang (3.4 a questo punto) e libclc. Presumo che tu abbia una conoscenza di base su come configurare e compilare LLVM e Clang, quindi ho appena elencato i flag di configurazione che ho usato.

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. compilazione LLVM e clang con supporto nvptx:

   ../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) e compilarlo.

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

Se avete problema compilazione di questo potrebbe essere necessario correggere paio di intestazioni in ./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. compilazione del kernel di LLVM IR assember :

   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. Collegamento kernel con implementazioni integrate da libclc

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

3. compilazione completamente legato 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

Con la versione corrente di llvm (3.4), libclc e nvptx back-end, il processo di compilazione è leggermente cambiato.

Devi dire esplicitamente al backend nvptx quale interfaccia driver usare; le opzioni disponibili sono nvptx-nvidia-cuda o nvptx-nvidia-nvcl (per OpenCL) e i loro equivalenti a 64 bit nvptx64-nvidia-cuda o nvptx64-nvidia-nvcl.

Il codice .ptx generato differisce leggermente in base all'interfaccia scelta. Nel codice assembly prodotto per l'API del driver CUDA, intrinsic .global e .ptr vengono rilasciati dalle funzioni di immissione ma sono richiesti da OpenCL.Ho modificato passi di compilazione di Mikael leggermente per produrre codice che possono essere eseguiti con OpenCL host:

1. Compila per 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. kernel Link:

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

3. Compile to Ptx:

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