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

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

Ответы

Ответ 1

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

Ответ 2

В текущей версии llvm (3.4), libclc и nvptx back-end процесс компиляции несколько изменился.

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

Сгенерированный код .ptx немного отличается в зависимости от выбранного интерфейса. В коде сборки, созданном для API драйвера CUDA, intrinsics.global и .ptr удаляются из функций ввода, но они требуются OpenCL. Я немного изменил шаги Mikael, чтобы создать код, который можно запустить с помощью хоста OpenCL:

Скомпилируйте 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

Ядро связи:

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

Скомпилируйте Ptx:

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

Ответ 3

Вот краткое руководство, как это сделать с стволом 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];
}

Скомпилируйте llvm и clang с поддержкой nvptx:

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

Получить 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"

Скомпилировать ядро для 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

Свяжите ядро со встроенными реализациями из libclc
```
llvm-link libclc/nvptx--nvidiacl/lib/builtins.bc square.ll -o square.linked.bc
```
Скомпилировать полностью привязанный 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;
    }