[llvm-bugs] [Bug 50081] New: Error in OpenCL kernel parameter: '__global [SOME_STRUCT] *__private' cannot be used as the type of a kernel parameter

[via llvm-bugs]

https://bugs.llvm.org/show_bug.cgi?id=50081

            Bug ID: 50081
           Summary: Error in OpenCL kernel parameter: '__global
                    [SOME_STRUCT] *__private' cannot be used as the type
                    of a kernel parameter
           Product: clang
           Version: trunk
          Hardware: PC
                OS: Linux
            Status: NEW
          Severity: enhancement
          Priority: P
         Component: OpenCL
          Assignee: unassignedclangbugs at nondot.org
          Reporter: drohr at jwdt.org
                CC: anastasia.stulova at arm.com, llvm-bugs at lists.llvm.org

Created attachment 24791
  --> https://bugs.llvm.org/attachment.cgi?id=24791&action=edit
testcase

The attached testcase fails to compile for me with C++ for OpenCL using clang
trunk (clang version 13.0.0 (https://github.com/llvm/llvm-project.git
1d96107cfec5539b02b4464009a6c330c0593e44)) giving several errors of the form:

test2.cl:30328:93: error: '__global o2::gpu::GPUConstantMem *__private' cannot
be used as the type of a kernel parameter
__kernel void krnl_GPUTPCNeighboursFinder(__global char *gpu_mem, __global
GPUConstantMem * pConstant, int iSlice_internal ) { __local typename
GPUTPCNeighboursFinder::GPUSharedMemory smem; GPUTPCNeighboursFinder::template
Thread<GPUTPCNeighboursFinder::defaultKernel>(get_num_groups(0),
get_local_size(0), get_group_id(0), get_local_id(0), smem,
GPUTPCNeighboursFinder::Processor((*pConstant))[iSlice_internal] ); }
                                                                               
            ^
test2.cl:30329:94: error: '__global o2::gpu::GPUConstantMem *__private' cannot
be used as the type of a kernel parameter
__kernel void krnl_GPUTPCNeighboursCleaner(__global char *gpu_mem, __global
GPUConstantMem * pConstant, int iSlice_internal ) { __local typename
GPUTPCNeighboursCleaner::GPUSharedMemory smem;
GPUTPCNeighboursCleaner::template
Thread<GPUTPCNeighboursCleaner::defaultKernel>(get_num_groups(0),
get_local_size(0), get_group_id(0), get_local_id(0), smem,
GPUTPCNeighboursCleaner::Processor((*pConstant))[iSlice_internal] ); }

Command to reproduce:
clang++ -cl-std=clc++ -x cl -emit-llvm --target=spir64-unknown-unknown -Xclang
-fdenormal-fp-math-f32=ieee -cl-mad-enable -cl-no-signed-zeros
-ferror-limit=1000 -Xclang -finclude-default-header
-Dcl_clang_storage_class_specifiers -c test2.cl -o test.bc

-- 
You are receiving this mail because:
You are on the CC list for the bug.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-bugs/attachments/20210422/384ca182/attachment.html>