[LLVMdev] Error at backend compilation when created a .ptx file for opencl use

[L Dumas]

Hi,

I 'm kind of new to llvm/clang/llc and when I tried to compile a simple
opencl code... well it kind of failed... 

Now it might be just me not doing what is required for the cross-compilation
to succeed, but my opencl code is not tricky though llvm fails giving me a
"Cannot select: 0x1c8a340: i32 = ExternalSymbol'log2f'" error. 

I do have a log function called in my code which basically does the
following : 

test.cl: 
float random 
( 
    float* ctr, 
    float* key 
) 
{ 
    *key = *ctr-1; 
    float b = *ctr; 
    float c = *key; 
    return (b+c)/(b+c+2.0f); 
} 
__kernel void kernel 
( 
    __global    float           *tab 
) 
{ 
    unsigned int idx = get_local_id(1) + get_local_size(1) *
(get_local_id(0) + get_local_size(0) * ( get_group_id(1) + get_group_id(0) *
get_num_groups(1))); 
    float configThr = (float)idx; 
    float etatThr = (float)idx; 
    tab[idx] = tab[idx] * log(1.f-random(&etatThr, &configThr)); 
}; 



As you'll notice this code makes no sense at all but that is not the point
as I tried to highlight how simple it is. 

I'm using the following lines to build a .ptx file : 
clang   \ 
        -target nvptx \ 
        -S \ 
        -std=CL1.1 \ 
        test.cl -o test.ll \ 
        -I/usr/local/include/clc/ \ 
        -include clc/clc.h \ 
        -Dcl_clang_storage_class_specifiers \ 
        -fmacro-backtrace-limit=0 \ 
        -emit-llvm 

llc -march nvptx  test.ll -mcpu=sm_20 -mattr=ptx30 -o test.ptx 

I have llvm/clang 3.4 from january 2014. 

Is this something I'm doing wrong ? 
Can you guys think of workaround I could try ? 

Anyway thank you for reading that, any hint will be greatly appreciated... 


--- 
Apparently when the variable returned by my random function can clearly be
identified as positive and lesser than 1 then it compiles just fine... 



--
View this message in context: http://llvm.1065342.n5.nabble.com/Error-at-backend-compilation-when-created-a-ptx-file-for-opencl-use-tp66872.html
Sent from the LLVM - Dev mailing list archive at Nabble.com.