Estern non risolto durante la compilazione di OpenCL in PTX utilizzando Clang?

Question

Sto seguendo le istruzioni a this SO answer ma quando provo a eseguire il file risultante PTX ottengo l'errore seguito in clBuild

ptxas fatal : Unresolved extern function 'get_group_id' 

Nel file PTX ho il seguente per ogni chiamata di funzione OpenCL Io uso

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

Quanto sopra non è presente nei file PTX creati dal runtime OpenCL quando fornisco un file CL. Invece ha il registro speciale appropriato.

seguito these instructions (link nei confronti di una biblioteca libclc diversa) mi dà un errore di segmentazione durante la LLVM IR per la compilazione PTX con il seguente errore:

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

sono quelle istruzioni ancora valido? C'è qualcos'altro che dovrei fare?

sto utilizzando l'ultima versione di libclc, Clang 3.7, e il driver Nvidia 352,39

Answer 1

Il problema è che LLVM non fornisce una libreria di codice del dispositivo OpenCL. llvm fornisce comunque gli elementi intrinseci per ottenere gli ID di un thread GPU. Ora devi scrivere i tuoi impianti di get_global_id ecc. Usando i builtin di clang e compilarlo in llvm bitcode con il target nvptx. Prima di abbassare il tuo IR su PTX usi llvm-link per collegare la libreria di dispositivi con il modulo OpenCL compilato e il gioco è fatto.

Un esempio di come si può scrivere una tale funzione:

#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; 
    } 
}