Как использовать clang для компиляции кода OpenCL в ptx?

Question

Clang 3.0 способен скомпилировать OpenCL в ptx и использовать инструмент Nvidia для запуска кода ptx на графическом процессоре. Как я могу это сделать? Пожалуйста, будьте конкретны.

Answer 1

См. Justin Holewinski's blog для конкретного примера или this thread для более подробных шагов и ссылок на образцы.

Answer 2

Вот краткое руководство, как это сделать с стволом Clang (3.4 на данный момент) и libclc. Я предполагаю, что у вас есть базовые знания по настройке и компиляции LLVM и Clang, поэтому я просто перечислил используемые флажки configure.

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. Compile LLVM и лязг с nvptx поддержки:

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

2. Получить libclc (GIT клон http://llvm.org/git/libclc.git) и скомпилировать его.

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

Если у вас есть проблемы с компиляцией этого вам, возможно, потребуется, чтобы исправить несколько заголовков в ./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. Compile ядра для 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. Ссылка на ядро ​​с встроенные реализации от libclc

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

3. Вкомпилировать полностью связаны 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

С текущей версией из LLVM (3.4), libclc и nvptx, процесс компиляции несколько изменился.

Вы должны явно указать бэкэнд nvptx, какой интерфейс драйвера использовать; ваши варианты: nvptx-nvidia-cuda или nvptx-nvidia-nvcl (для OpenCL) и их 64-разрядные эквиваленты nvptx64-nvidia-cuda или nvptx64-nvidia-nvcl.

Сгенерированный код .ptx немного отличается в зависимости от выбранного интерфейса. В коде сборки, созданном для API драйвера CUDA, intrinsics .global и .ptr удаляются из функций ввода, но они требуются OpenCL.Я изменил компиляции шаги Микаэл слегка для производства кода, которые могут выполняться с 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. Скомпилировать до Ptx:

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