Crash on Dynamic Parallelism in OpenCL 2.x

[mshockwave]

Hi, I just want to compiler a simple vector add example(adapted from NVIDIA's example) with dynamic parallelism in OpenCL 2.x but end up getting crashed.
Here is the CL kernel code:

__kernel void VectorAdd(__global const float* a, __global const float* b, __global float* c, int iNumElements)
{
    int tid = get_global_id(0);
    device_queue devQ = get_default_queue();
    ndrange ndrange1(32);

    auto myblock = [=](void)->void{
      int iGID = get_global_id(0) + tid * 32;
      // add the vector elements
      c[iGID] = a[iGID] + b[iGID];
    };

    auto err_ret = devQ.enqueue_kernel(enqueue_policy::wait_kernel,
                                      ndrange1,
                                      myblock);

}

The error I got is:

clang: /path/to/llvm/lib/SPIRV/SPIRVWriter.cpp:464: SPIRV::SPIRVType* SPIRV::LLVMToSPIRV::transType(llvm::Type*): Assertion `!ET->isFunctionTy() && "Function pointer type is not allowed"' failed.

After I turn on the SPIRV debugging message, I got more detailed information about the location of the error:

[mapValue]   call spir_func void @"_ZNU3AS42cl9__details16__enqueue_helperIZ9VectorAddE3$_0NS0_8__paramsIJEEENS0_5__seqImJEEES6_S4_S4_E32__get_enqueue_kernel_static_dataEOU3AS4S2_OU3AS4NS_5tupleIJEEE"(%class.anon.0 addrspace(4)* sret %1, %class.anon addrspace(4)* dereferenceable(32) %0, %"class.cl::tuple<>" addrspace(4)* dereferenceable(1) %args) => 99�����
[transValue]   %call2 = call spir_func void (i8 addrspace(4)*)* ()* @"_ZNU3AS42cl9__details16__enqueue_helperIZ9VectorAddE3$_0NS0_8__paramsIJEEENS0_5__seqImJEEES6_S4_S4_E28__get_enqueue_kernel_wrapperEv"()
[transType] void (i8 addrspace(4)*)* ()
[transType] void (i8 addrspace(4)*)*

It seems that the problem is raised in libclcxx. When the SPIRV backend use transType(llvm::Type*) to transform function parameters, transType(llvm::Type*) seem to assert that if the input Type is pointer type, then it must not be function pointer type. But it's a little bit strange since there are many chances in OpenCL 2.x that function parameters are lambda functions or function pointers.

[bsochack]

Thank you for reporting this issue. I know how to fix it, please give me a few days.

[mshockwave]

Cheers! I'm looking forward.

[bsochack]

Please update libclcxx to the latest version. I believe it should fix your issue

[mshockwave]

I works! Thanks for rapid fix :)