¿Cómo usar clang para compilar código OpenCL a ptx?

Question

Clang 3.0 puede compilar OpenCL a ptx y usar la herramienta de Nvidia para iniciar el código ptx en la GPU. ¿Cómo puedo hacer esto? Por favor sea especifico.

Answer 1

Consulte Justin Holewinski's blog para obtener un ejemplo específico o this thread para obtener algunos pasos más detallados y vínculos a las muestras.

Answer 2

Aquí hay una breve guía de cómo hacerlo con Clang trunk (3.4 en este punto) y libclc. Supongo que tiene conocimientos básicos sobre cómo configurar y compilar LLVM y Clang, por lo que acabo de enumerar los indicadores de configuración que he utilizado.

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 Compilar y sonido metálico con soporte nvptx:

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

2. Obtener libclc (git clone http://llvm.org/git/libclc.git) y compilarlo.

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

Si tiene problemas para compilar este puede que tenga que corregir par de cabeceras en ./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. de compilación del kernel a Assember 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. Enlace kernel con implementaciones incorporadas de libclc

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

3. Compilar totalmente vinculado LLVM IR a 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

Con la versión actual de de llvm (3.4), back-end libclc y nvptx, el proceso de compilación ha cambiado ligeramente.

Tiene que decirle explícitamente al back-end nvptx qué interfaz del controlador usar; sus opciones son nvptx-nvidia-cuda o nvptx-nvidia-nvcl (para OpenCL) y sus equivalentes de 64 bits nvptx64-nvidia-cuda o nvptx64-nvidia-nvcl.

El código .ptx generado difiere ligeramente según la interfaz elegida. En el código ensamblado producido para la API del controlador CUDA, los intrínsecos .global y .ptr se eliminan de las funciones de entrada, pero son requeridos por OpenCL.He modificado los pasos de compilación de Mikael ligeramente para producir código que se pueden ejecutar con OpenCL anfitrión:

1. Compilar a 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. Enlace núcleo:

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

3. Compilar a Ptx:

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