Cl get kernel arg info

[franz]

Please review, thanks

[franz]

I have rebased it onto pull request 92, and still need to add a SPIR variant of the test

[kraiskil]

Looks good to me & passes tests at this point. But won't merge if there is more coming?

[franz]

I enhanced the test code to also test SPIR bitcode, and made the code a bit more robust; i think all of the TODOs for this branch are done, so it should be ready.

[pjaaskel]

OK, I'll merge if tests pass.

[pjaaskel]

Putting this here also. Most of the previous tests now fail with:

llvm-3.4/include/llvm/Support/Casting.h:239: typename llvm::cast_retty<X, Y*>::ret_type llvm::cast(Y*) [with X = llvm::MDString; Y = llvm::Value; typename 
                 llvm::cast_retty<X, Y*>::ret_type = llvm::MDString*]: Assertion `isa<X>(Val) && "cast<Ty>() argument of incompatible type!"' failed.

[pjaaskel]

lt-example1: /home/visit0r/local/stow/llvm-3.4/include/llvm/Support/Casting.h:239: typename llvm::cast_retty<X, Y*>::ret_type llvm::cast(Y*) [with X = llvm::MDString; Y = llvm::Value; typename llvm::cast_retty<X, Y*>::ret_type = llvm::MDString*]: Assertion `isa<X>(Val) && "cast<Ty>() argument of incompatible type!"' failed.

Program received signal SIGABRT, Aborted.
0x00007ffff745e475 in *__GI_raise (sig=<optimized out>) at ../nptl/sysdeps/unix/sysv/linux/raise.c:64
64      ../nptl/sysdeps/unix/sysv/linux/raise.c: No such file or directory.
(gdb) bt
#0  0x00007ffff745e475 in *__GI_raise (sig=<optimized out>) at ../nptl/sysdeps/unix/sysv/linux/raise.c:64
#1  0x00007ffff74616f0 in *__GI_abort () at abort.c:92
#2  0x00007ffff7457621 in *__GI___assert_fail (assertion=0x7ffff670c368 "isa<X>(Val) && \"cast<Ty>() argument of incompatible type!\"", file=<optimized out>, line=239, 
    function=0x7ffff670d640 "typename llvm::cast_retty<X, Y*>::ret_type llvm::cast(Y*) [with X = llvm::MDString; Y = llvm::Value; typename llvm::cast_retty<X, Y*>::ret_type = llvm::MDString*]") at assert.c:81
#3  0x00007ffff5c8693a in llvm::cast<llvm::MDString, llvm::Value> (Val=<optimized out>) at /home/visit0r/local/stow/llvm-3.4/include/llvm/Support/Casting.h:239
#4  0x00007ffff5cfe20d in _M_dispose (__a=..., this=<optimized out>) at /usr/include/c++/4.7/bits/basic_string.h:246
#5  ~basic_string (this=<optimized out>, __in_chrg=<optimized out>) at /usr/include/c++/4.7/bits/basic_string.h:536
#6  pocl_llvm_get_kernel_arg_metadata (kernel_name=kernel_name@entry=0x4017dd "dot_product", input=input@entry=0x6c1d50, kernel=kernel@entry=0x60af30) at pocl_llvm_api.cc:393
#7  0x00007ffff5cff80a in pocl_llvm_get_kernel_metadata (program=program@entry=0x608d30, kernel=kernel@entry=0x60af30, device_i=device_i@entry=0, kernel_name=kernel_name@entry=0x4017dd "dot_product", 
    device_tmpdir=device_tmpdir@entry=0x7fffffffcaa0 "/tmp/poclv3jz96/pthread", descriptor_filename=descriptor_filename@entry=0x7fffffffcea0 "/tmp/poclv3jz96/pthread/dot_product/descriptor.so", errcode=errcode@entry=0x7fffffffca9c)
    at pocl_llvm_api.cc:741
#8  0x00007ffff5cec31b in POclCreateKernel (program=0x608d30, kernel_name=0x4017dd "dot_product", errcode_ret=0x0) at clCreateKernel.c:88
#9  0x00007ffff79bd69f in clCreateKernel () from /usr/lib/libOpenCL.so.1
#10 0x00000000004013c8 in exec_dot_product_kernel (
    program_source=program_source@entry=0x602250 "#define USE_VECTOR_DATATYPES\n\n__kernel void \ndot_product (__global const float4 *a,  \n\t     __global const float4 *b, __global float *c) \n{ \n  int gid = get_global_id(0); \n\n#ifndef USE_VECTOR_DATATYPE"..., n=n@entry=128, srcA=srcA@entry=0x602630, srcB=srcB@entry=0x602e40, dst=dst@entry=0x602010) at example1_exec.c:109
#11 0x0000000000400f82 in main () at example1.c:79

Debian Stable, LLVM 3.4.1, x86_64.

It seems valid: llvm::MDString *kernel_prototype = llvm::cast(kernel_iter->getOperand(0)); You are casting llvm::Function to llvm::MDString. I wonder how it works for you but not for me. See Workgroup.cc:634 -- it's the same MD there, right?

[franz]

I changed the cast to llvm::Function. I have no idea how could it work before..

[pjaaskel]

OK, that cast crash is now gone, and other tests have passed so far, but the new test failed:

./testsuite-runtime.at:67: $abs_top_builddir/tests/runtime/test_clGetKernelArgInfo
stderr:
%opencl.image2d_t.1*lt-test_clGetKernelArgInfo: pocl_llvm_api.cc:645: int pocl_llvm_get_kernel_metadata(cl_program, cl_kernel, int, const char *, const char *, char *, int *): Assertion `p->getAddressSpace() == 2' failed.
/home/visit0r/src/pocl-public/tests/testsuite.dir/at-groups/97/test-source: line 24:  7061 Aborted                 $abs_top_builddir/tests/runtime/test_clGetKernelArgInfo
stdout:

NON-SPIR

SPIR with metadata
opening File: /home/visit0r/src/pocl-public/tests/runtime/clGetKernelArgInfo.spir64_meta
program size: 2596

[pjaaskel]

The issue is exposing another issue with the binary loading interface. The (old) problem of linking together LLVM::Modules with the same global type names (in this case %opencl.image2d_t) appears to be not fixed when having both source kernels and binary (SPIR) kernels in the same program.

When LLVM bitcode linker links together two Modules that use different LLVMContext and have the same named type, it creates a new type %opencl.image2d_t.N with incremented N in order to avoid accidentally wrongly mergins types (it cannot tell if the other Module really meant the same type). This causes the check that figures out if the parameter is a image to fail. I do not think we can just extend the check to match opencl.image2d_t.* because that caused some other problems which I cannot remember (at least when we actually introduce the struct that replaces this opaque type hand there might be some problems).

Thus, the way this has been working is to use the same singleton LLVMContext for all bitcode (Module) operations. LLVMContext keeps book of the types so when we introduce two Modules in the same LLVMContext that use the same type (name), they should merge to the same type. For some reason this now fails, which probably means that the same LLVMContext is not propagated to some IR loading functions. You can see that it works if you only load from the SPIR in that test, but not if you first load from a source and then from SPIR.

Unfortunately I do not have now time to debug this further, but because this is a bug unrelated to the feature at hand, I suggest you somehow split the test to two (one for src and one for SPIR) to circumvent the issue and still test both cases so we get your work merged. I opened issue #95 for the actual problem. Of course, if you feel like digging into the LLVM and pocl more, you could try to figure it out yourself, because (due to holidays) I do not know when I can get into it.