ओपनसीएल को पीटीएक्स कोड में संकलित करने के लिए क्लैंग का उपयोग कैसे करें?

Question

क्लैंग 3.0 ओपनसीएल को पीटीएक्स में संकलित करने और जीपीयू पर पीटीएक्स कोड लॉन्च करने के लिए एनवीडिया के टूल का उपयोग करने में सक्षम है। मैं यह कैसे कर सकता हूँ? कृपया विशिष्ट रहें।

Answer 1

कुछ विशिष्ट विस्तृत चरणों और नमूने के लिंक के लिए Justin Holewinski's blog एक विशिष्ट उदाहरण या this thread देखें।

Answer 2

यहां संक्षिप्त मार्गदर्शिका है कि क्लेंग ट्रंक (3.4 इस बिंदु पर) और libclc के साथ इसे कैसे करें। मुझे लगता है कि आपके पास एलएलवीएम और क्लैंग को कॉन्फ़िगर और संकलित करने के लिए बुनियादी ज्ञान है, इसलिए मैंने अभी कॉन्फ़िगर किए गए कॉन्फ़िगरेशन फ़्लैग सूचीबद्ध किए हैं।

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 
   

आप समस्या यह संकलन है, तो आप LLVM आईआर Assember को ./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. संकलित कर्नेल में हेडर की जोड़ी ठीक करने के लिए आवश्यकता हो सकती है :

   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. संकलित पूरी तरह से PTX

   को LLVM आईआर जुड़ा हुआ

   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 (ओपनसीएल के लिए) और उनके 64 बिट समकक्ष nvptx64-nvidia-cuda या nvptx64-nvidia-nvcl हैं।

जेनरेट किया गया .ptx कोड चयनित इंटरफ़ेस के अनुसार थोड़ा अलग है। CUDA ड्राइवर API के लिए उत्पादित असेंबली कोड में, इंट्रिनिक्स .global और .ptr को एंट्री फ़ंक्शंस से हटा दिया जाता है लेकिन उन्हें ओपनसीएल द्वारा आवश्यक होता है।

1. संकलित LLVM के लिए आईआर:

   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. मैं माइकल के संकलन चरण कुछ कोड का उत्पादन करने कि OpenCL मेजबान के साथ चलाया जा सकता है संशोधित कर लिया है पीटीएक्स के लिए संकलित:

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