使用 Clang 将 OpenCL 编译为 PTX 时未解决的外部问题？答案

【问题标题】：Unresolved extern when compiling OpenCL to PTX using Clang?使用 Clang 将 OpenCL 编译为 PTX 时未解决的外部问题？
【发布时间】：2015-12-03 09:54:52
【问题描述】：

我正在按照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
)
;

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

按照these instructions（链接到不同的 libclc 库）在 LLVM IR 到 PTX 编译期间给我一个分段错误，并出现以下错误：

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

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

我正在使用最新版本的 libclc、Clang 3.7 和 Nvidia 驱动程序 352.39

【问题讨论】：

标签： clang opencl llvm-clang

【解决方案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; 
  } 
}

【讨论】：