Làm thế nào để sử dụng clang để biên dịch OpenCL thành mã ptx?

Question

Clang 3.0 có thể biên dịch OpenCL thành ptx và sử dụng công cụ của Nvidia để khởi chạy mã ptx trên GPU. Tôi có thể làm cái này như thế nào? Hãy cụ thể.

Answer 1

Xem Justin Holewinski's blog để biết ví dụ cụ thể hoặc this thread để biết thêm một số bước chi tiết và liên kết đến mẫu.

Answer 2

Dưới đây là hướng dẫn ngắn gọn cách thực hiện với thân cây Clang (3.4 tại điểm này) và libclc. Tôi cho rằng bạn có kiến ​​thức cơ bản về cách cấu hình và biên dịch LLVM và Clang, vì vậy tôi chỉ liệt kê các cờ cấu hình mà tôi đã sử dụng.

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 và kêu vang với sự hỗ trợ nvptx:

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

2. Nhận libclc (git clone http://llvm.org/git/libclc.git) và biên dịch nó.

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

Nếu bạn có vấn đề biên dịch này, bạn có thể cần phải sửa chữa vài tiêu đề trong ./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. kernel 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. Liên kết nhân với triển khai xây dựng trong từ libclc

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

3. Compile liên kết đầy đủ 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

Với phiên bản hiện tại của các llvm (3.4), libclc và nvptx back-end, quá trình biên dịch đã thay đổi một chút.

Bạn phải nói rõ với phần phụ trợ nvptx mà giao diện trình điều khiển sử dụng; các tùy chọn của bạn là nvptx-nvidia-cuda hoặc nvptx-nvidia-nvcl (cho OpenCL) và 64 bit tương đương nvptx64-nvidia-cuda hoặc nvptx64-nvidia-nvcl của chúng.

Mã .ptx được tạo khác nhau tùy theo giao diện đã chọn. Trong mã assembly được tạo ra cho API trình điều khiển CUDA, nội tại .global và .ptr bị loại bỏ khỏi các hàm mục nhập nhưng chúng được yêu cầu bởi OpenCL.Tôi đã sửa đổi các bước biên dịch Mikael của hơi để sản xuất mã mà có thể chạy với OpenCL host:

1. Compile để 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. liên kết hạt nhân:

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

3. Biên dịch sang Ptx:

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