如何使用clang編譯OpenCL到ptx代碼？

Question

Clang 3.0能夠將OpenCL編譯爲ptx，並使用Nvidia的工具在GPU上啓動ptx代碼。我怎樣才能做到這一點？請具體說明。

Answer 1

查看​​的具體示例或this thread的一些更詳細的步驟和示例鏈接。

Answer 2

這裏是簡單的指南如何做到這一點叮噹幹線（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編譯和鐺與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. 編譯內核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

隨着LLVM的當前版本（ 3.4），libclc和nvptx後端，編譯過程略有改變。

您必須明確告訴nvptx後端使用哪個驅動程序接口;你的選擇是nvptx-nvidia-cuda或nvptx-nvidia-nvcl（適用於OpenCL）及其64位相當於nvptx64-nvidia-cuda或nvptx64-nvidia-nvcl。

生成的.ptx代碼根據所選的接口略有不同。在爲CUDA驅動程序API生成的彙編代碼中，intrinsics .global和.ptr從輸入函數中刪除，但它們是OpenCL需要的。我修改的Mikael的小幅編譯步驟產生的代碼可與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