kernel compiler crash with a private constant array

[Pekka Jääskeläinen]

"void test_arrays_float(int const *const data, int const ndata)\n"
"{\n"
" float s = 0.0f;\n"
" for (int i=0; i<ndata; ++i) {\n"
" s += as_float(data[i]);\n"
" }\n"
" printf(\"sum=%g\\n\", s);\n"
" for (;;);\n"
"}\n"
"\n"
"kernel void test_kernel(global int *input, global int* output)\n"
"{\n"
" int const data[] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 11};\n"
" int const ndata = sizeof(data) / sizeof(*data);\n"
" test_arrays_float(data, ndata);\n"
"}\n";

This one causes a crash when copying the call to the test_arrays_float, related somehow to the data array being passed to the test_arrays_float. Didn't debug it much yet.

[Pekka Jääskeläinen (pekka-jaaskelainen)]

It chokes on a kernel like this:

define void @test_kernel(i32 addrspace(3)* nocapture %input, i32 addrspace(3)* nocapture %output, [11 x i32]* %_local0) noreturn nounwind uwtable {
entry:
  tail call void @test_arrays_float(i32* getelementptr inbounds ([11 x i32]* @test_kernel.data, i64 0, i64 0), i32 11)
  unreachable
}

opt: /home/visit0r/src/llvm-3.1.src/include/llvm/Support/Casting.h:194: typename llvm::cast_retty<To, From>::ret_type llvm::cast(const Y&) [with X = llvm::Constant, Y = llvm::Value*]: Assertion `isa<X>(Val) && "cast<Ty>() argument of incompatible type!"' failed.

#0 0x00007ffff62a9b9e in __libc_waitpid (pid=<optimized out>, stat_loc=0x7fffffffc030, options=<optimized out>) at ../sysdeps/unix/sysv/linux/waitpid.c:32
#1 0x00007ffff7335b5a in llvm::sys::Program::Wait (this=0x7fffffffc0a0, path=..., secondsToWait=0, ErrMsg=0x7fffffffc160) at Unix/Program.inc:326
#2 0x00007ffff7337417 in llvm::sys::Program::ExecuteAndWait (path=..., args=<optimized out>, envp=0x0, redirects=<optimized out>, secondsToWait=0, memoryLimit=<optimized out>, ErrMsg=0x7fffffffc160) at Program.cpp:35
#3 0x00007ffff7322cf1 in ExecGraphViewer (ErrMsg=<optimized out>, wait=<optimized out>, Filename=<optimized out>, args=<optimized out>, ExecPath=<optimized out>) at GraphWriter.cpp:61
#4 llvm::DisplayGraph (Filename=..., wait=true, program=<optimized out>) at GraphWriter.cpp:196
#5 0x00007ffff678bbef in ViewGraph<llvm::Function const*> (Title=<optimized out>, ShortNames=<optimized out>, Name=<optimized out>, G=<optimized out>, Program=<optimized out>) at /home/visit0r/src/llvm-3.1.src/include/llvm/Support/GraphWriter.h:352
#6 llvm::Function::viewCFG (this=<optimized out>) at CFGPrinter.cpp:146
#7 <function called from gdb>
#8 0x00007ffff55bb1b5 in *__GI_raise (sig=<optimized out>) at ../nptl/sysdeps/unix/sysv/linux/raise.c:64
#9 0x00007ffff55bdfc0 in *__GI_abort () at abort.c:92
#10 0x00007ffff55b4301 in *__GI___assert_fail (assertion=0x7ffff74ebd50 "isa<X>(Val) && \"cast<Ty>() argument of incompatible type!\"", file=<optimized out>, line=194,
    function=0x7ffff74f5120 "typename llvm::cast_retty<To, From>::ret_type llvm::cast(const Y&) [with X = llvm::Constant, Y = llvm::Value*]") at assert.c:81
#11 0x00007ffff6796b7f in llvm::cast<llvm::Constant, llvm::Value*> (Val=<optimized out>) at /home/visit0r/src/llvm-3.1.src/include/llvm/Support/Casting.h:194
#12 0x00007ffff73c77bf in llvm::MapValue (V=0x630758, VM=..., Flags=llvm::RF_NoModuleLevelChanges, TypeMapper=0x0) at ValueMapper.cpp:132
#13 0x00007ffff73c7ecb in llvm::RemapInstruction (I=0x6313b8, VMap=..., Flags=llvm::RF_NoModuleLevelChanges, TypeMapper=0x0) at ValueMapper.cpp:164
#14 0x00007ffff7379639 in llvm::CloneFunctionInto (NewFunc=<optimized out>, OldFunc=<optimized out>, VMap=..., ModuleLevelChanges=<optimized out>, Returns=<optimized out>, NameSuffix=<optimized out>, CodeInfo=0x0, TypeMapper=0x0) at CloneFunction.cpp:148
#15 0x00007ffff5357437 in (anonymous namespace)::GenerateHeader::ProcessAutomaticLocals (this=<optimized out>, F=0x633f90, out=<optimized out>) at GenerateHeader.cc:292
...

The assert happens in llvm::CloneFunctionInto. The curious thing is the first argument to the...

[Pekka Jääskeläinen (pekka-jaaskelainen)]

mri-gridding of Parboil hits this also therefore I investicated the issue.

It's because of automatic local handling. Clang converts automatic __locals to global variables. These have to be converted to kernel function arguments so it's possible to task parallelize multiple work group executions (so they can use their separate local buffers). However, as it seems like a constant global buffer to the LLVM, it generates "constant GEPs" which in this case trigger the crash when replicating the kernel function. These constant GEP references to the local variables cannot be converted to references to the generated local function arguments.

Therefore, I think the fix is to convert the constant GEPs referring to the temporary local variables to regular values before trying to copy the kernel.

[Pekka Jääskeläinen (pekka-jaaskelainen)]

int test(__local int* x)
{
    return *x++;
}

__kernel void foo(void) {
    __local int i[2];
    i[0] = test(&i[0]);
}

A minimized test case that produces the constant GEP to a function call argument list.

[Pekka Jääskeläinen (pekka-jaaskelainen)]

Converted them to normal GEPs using a pass from the SAFECode project.