Displaying 2 results from an estimated 2 matches for "if_eq".
Did you mean:
if_ed
2013 Dec 31
4
[LLVMdev] [Patch][RFC] Change R600 data layout
...it arithmetic except for load/store
operations that roughly map onto getelementptr.
The current datalayout for r600 includes n32:64, which is odd
because r600 cannot actually do any 64bit arith natively. This causes
particular problems in the optimizer with the following kernel:
__kernel void if_eq(__global unsigned long* out, unsigned arg0)
{
int i=0;
for(i = 0; i < arg0; i++){
out[i] = i;
}
}
Clang decides that instead of adding a sext i32 %i to i64 before
getelementptr, it would be best to just go ahead and promote the
variable i to i64. Which would be all...
2013 Dec 31
2
[LLVMdev] [PATCH] R600 - Fix zero extend of i1
Hi,
When trying to compile a trivial opencl kernel such as:
__kernel void if_eq(__global int * out, int arg0, int arg1){
out[0] = arg0==arg1?0:1;
}
Clang generates IR like:
%1 = icmp eq i32 %arg0, %arg1
%. = zext i1 %1 to i32
This eventually crashes ISel on R600. Attached patch adds a selector so
it will compile.
Regards,
Jon Pry
jonpry at gmail.com
--------...