search for: get_work_dim

...the compilation unit. In OpenCL the user program can invoke a kernel as an NDRange, i.e. with an implied loop around it to iterate over an index space of 1 to 3 dimensions. (This is the "big idea" of OpenCL). (The index values are available in the function body from intrinsic functions get_work_dim() and get_global_id(uint workdim).) But that implied loop is only applied when directly called from the user program. When a kernel is called from another kernel, it behaves as a regular function call and just adopts the caller's index point. The spec does not specify whether or how that impl...

On 16 December 2010 14:33, Peter Collingbourne <peter at pcc.me.uk> wrote: > On Thu, Dec 16, 2010 at 06:16:25PM -0000, Anton Lokhmotov wrote: > > Mike Gist wrote: > > > You could also consider placing all kernel functions in a 'kernel' > > > section, or adding a function attribute for kernels. > > Unlike in Clang, function attribute bit-fields in

...> > declare float @llvm.sqrt.f32(float) nounwind readonly > > declare float @llvm.exp.f32(float) nounwind readonly > > declare float @llvm.log.f32(float) nounwind readonly > > declare float @fabs(float) > > define void @BAD(%1* noalias nocapture %arg_struct, i32 %get_work_dim, i32* %get_global_size, i32* %get_local_size, i32* %get_group_id) { > entry: > %0 = getelementptr %1* %arg_struct, i64 0, i32 0 > %1 = load float addrspace(1)** %0, align 8 > %2 = getelementptr %1* %arg_struct, i64 0, i32 2 > %3 = load float addrspace(1)** %2, align 8 > %4 =...

On Thu, Dec 16, 2010 at 06:16:25PM -0000, Anton Lokhmotov wrote: > Mike Gist wrote: > > You could also consider placing all kernel functions in a 'kernel' > > section, or adding a function attribute for kernels. > Unlike in Clang, function attribute bit-fields in LLVM are pretty crowded > (only couple of bits are unused?). Besides, we do not want to represent >

Hello everybody, I have stumbled upon a test case (the attached module is a slightly reduced version) that shows extremely reduced performance on linux compared to windows when executed using LLVM's JIT. We narrowed the problem down to the actual code being generated, the source IR on both systems is the same. Try compiling the attached module: llc -O3 -filetype=asm -o BAD.s BAD.ll Under