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 specificoCome usare clang per compilare OpenCL al codice ptx?
risposta
Vedere Justin Holewinski's blog per un esempio specifico o this thread per alcuni passaggi più dettagliati e collegamenti agli esempi.
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];
}
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
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"
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
Collegamento kernel con implementazioni integrate da libclc
llvm-link libclc/nvptx--nvidiacl/lib/builtins.bc square.ll -o square.linked.bc
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;
}
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:
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
kernel Link:
llvm-link libclc/built_libs/nvptx64--nvidiacl.bc test.ll -o test.linked.bc
Compile to Ptx:
clang -target nvptx64-nvidia-nvcl test.linked.bc -S -o test.nvptx.s
Per me ho dovuto passare i due ingressi al punto # 2 per farlo collegare correttamente. Fonte: https://groups.google.com/forum/#!msg/llvm-dev/Iv_u_3wh4lU/XINHv5HbAAAJ – Andrew
Il collegamento al blog non funziona più. Inoltre, se ricordo bene, si trattava di informazioni deprecate. –
Ho praticamente risolto il link del blog. – sschuberth