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

Question

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

Answer 1

С текущей версией бэкэнда llvm (3.4), libclc и nvptx процесс компиляции изменился незначительно.

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

Сгенерированный код .ptx немного отличается в зависимости отвыбранный интерфейс.В коде сборки, созданном для API драйвера CUDA, встроенные функции .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
   

Answer 2

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

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. Скомпилируйте llvm и clang с поддержкой nvptx:

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

2. Получите libclc (git clone 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. Скомпилировать ядро ​​в сборку 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. Ссылка на ядро ​​со встроенными реализациями из 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

См. блог Джастина Холевински для конкретного примера или эту ветку для некоторых более подробных шагов и ссылок на образцы.