如何使用clang将OpenCL编译为ptx代码？

Question

Clang 3.0能够将OpenCL编译为ptx,并使用Nvidia的工具在GPU上启动ptx代码.我怎样才能做到这一点？请具体说明.

Answer 1

使用当前版本的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
   

Answer 2

请参阅Justin Holewinski 的博客以获取特定示例或此线程以获取更详细的步骤和示例链接。

Answer 3

以下是如何使用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. 使用nvptx支持编译llvm和clang:

   ../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;
    }