使用Clang編譯OpenCL到PTX時無法解析的extern？

Question

我下面就this SO answer的說明，但是當我嘗試運行所產生的PTX文件，我得到了如下錯誤clBuild

ptxas fatal : Unresolved extern function 'get_group_id' 

在PTX文件每OpenCL的函數調用，我有以下我用

.func (.param .b64 func_retval0) get_group_id 
(
     .param .b32 get_group_id_param_0 
) 
; 

上面的內容在OpenCL運行時創建的PTX文件中不存在，當我將它提供給CL文件時。相反，它有適當的專用寄存器。

繼these instructions的LLVM IR期間，出現以下錯誤（針對不同的libclc庫鏈接）給了我一個分段故障到PTX編譯：

fatal error: error in backend: Cannot cast between two non-generic address spaces 

是這些指令是否仍然有效？還有什麼我應該做的？

我使用的是最新版本libclc，鐺3.7和NVIDIA驅動程序352.39

Answer 1

的問題是，LLVM不提供OpenCL設備代碼庫。然而，llvm提供了獲取GPU線程ID的內在機制。現在你必須使用clang的內建函數編寫自己的get_global_id等植入程序，並將它編譯爲具有nvptx目標的llvm位碼。在將IR降至PTX之前，您可以使用llvm-link將您的設備庫與編譯的OpenCL模塊鏈接，就是這樣。

一個例子，你會怎麼寫這樣的功能：

#define __ptx_mad(a,b,c) ((a)*(b)+(c)) 

__attribute__((always_inline)) unsigned int get_global_id(unsigned int dimindx) { 
    switch (dimindx) { 
    case 0: return __ptx_mad(__nvvm_read_ptx_sreg_ntid_x(), __nvvm_read_ptx_sreg_ctaid_x(), __nvvm_read_ptx_sreg_tid_x()); 
    case 1: return __ptx_mad(__nvvm_read_ptx_sreg_ntid_y(), __nvvm_read_ptx_sreg_ctaid_y(), __nvvm_read_ptx_sreg_tid_y()); 
    case 2: return __ptx_mad(__nvvm_read_ptx_sreg_ntid_z(), __nvvm_read_ptx_sreg_ctaid_z(), __nvvm_read_ptx_sreg_tid_z()); 
    default: return 0; 
    } 
}