Неразрешенный extern при компиляции OpenCL в PTX с помощью Clang?

Я следую инструкциям этого SO-ответа, но когда я пытаюсь запустить полученный файл PTX, я получаю следующую ошибку в clBuild

В файле PTX для каждого вызова функции OpenCL я использую следующее:

Вышеописанное не содержится в файлах PTX, созданных во время выполнения OpenCL, когда я предоставляю ему файл CL. Вместо этого он имеет собственный специальный регистр.

Следуя этим инструкциям (ссылки на другую библиотеку libclc) дает мне ошибку сегментации во время компиляции LLVM IR to PTX со следующей ошибкой:

Сохраняются ли эти инструкции? Есть ли что-то еще, что я должен делать?

Я использую последнюю версию libclc, Clang 3.7 и драйвер Nvidia 352.39

Ответы

Ответ 1

Проблема заключается в том, что llvm не предоставляет библиотеку кода устройства OpenCL. Однако llvm предоставляет встроенные средства для получения идентификаторов потока GPU. Теперь вам нужно написать свои собственные имплантации get_global_id и т.д., Используя встроенные встроенные clang и скомпилировать их с llvm-битом с целью nvptx. Перед тем, как опустить 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; 
  } 
}