llvm dev - Mar 2014 - [LLVMdev] Error at backend compilation when created a .ptx file for opencl use

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.

So I gave it a little more work...

Those lines are sufficient for the problem to appear :
__kernel void lancementKernel 
(
    __global    float           *tabPhotons
) 
{ 
    // idx est l'indice du thread considéré
    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)));
    tabPhotons[idx] = tabPhotons[idx] * log(1.f-(float)idx);
};


I have used -fmacro-backtrace-limit=0 on clang and the .i file generated
actually shows that my log call has been replaced by _clc_log2 which leads
me to llvm.log2 for float type.

It looks like there could be no implementation of log2 for float type. Does
that ring a bell to someone ? I have no clue about where the _clc_functions
are supposed to be defined. I feel like this should not be a hard issue but
I'm still trying to understand how clang/llvm works since I'm not
familiar
with the whole cross-compilation process.


I hope I am not wasting your time with some misunderstanding of my own....
Thank you for taking the time to read this post, hopefully you will know how
to help me get through that situation.



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