使用Clang将OpenCL编译为PTX时未解决的extern？

Question

我按照这个SO答案的说明但是当我尝试运行生成的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
)
;

当我提供一个CL文件时,OpenCL运行时创建的PTX文件中不存在上述内容.相反,它有适当的特殊寄存器.

遵循这些说明(链接到不同的libclc库)在LLVM IR到PTX编译期间给出了分段错误,并出现以下错误:

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

这些说明仍然有效吗？还有什么我应该做的吗？

我使用的是最新版本的libclc,Clang 3.7和Nvidia驱动程序352.39

Answer 1

问题是llvm不提供OpenCL设备代码库。然而，llvm 提供了用于获取 GPU 线程 ID 的内在函数。现在，您必须get_global_id使用 clang 的内置函数编写自己的植入等，并使用 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; 
  } 
}