Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[OpenCL] __opencl_c_work_group_collective_functions is not defined for SPIR-V target on OpenCL 3.0 #55770

Closed
haonanya opened this issue May 30, 2022 · 3 comments
Labels

Comments

@haonanya
Copy link

haonanya commented May 30, 2022

Hi, @AnastasiaStulova, @svenvh, @azabazno
__opencl_c_work_group_collective_functions is not defined for SPIRV https://github.com/llvm/llvm-project/blob/release/14.x/clang/lib/Headers/opencl-c-base.h#L67,L75, the test has compiling errors with the command clang+llvm-14.0.0-x86_64-linux-gnu-ubuntu-18.04/bin/clang++ -cc1 -include opencl-c.h -cl-std=CL3.0 -x cl -O2 -emit-llvm-bc -triple spir64 work_group_all.cl :
work_group_all.cl:5:18: error: implicit declaration of function 'work_group_all' is invalid in OpenCL
int result = work_group_all((input[tid] > input[tid+1]));
^
./work_group_all.cl:5:18: note: did you mean 'sub_group_all'?
./opencl-c.h:16534:23: note: 'sub_group_all' declared here
int __ovld __conv sub_group_all(int predicate);

Here is work_group_all.cl source:
__kernel void test_wg_all(global float *input, global int *output)
{
int tid = get_global_id(0);

int result = work_group_all((input[tid] > input[tid+1]));
output[tid] = result;

}

The test compiles successfully when add -D__opencl_c_work_group_collective_functions=1 or add __opencl_c_work_group_collective_functions definition for SPIRV on opencl-c-base.h.
Is the __opencl_c_work_group_collective_functions missing on opencl-c-base.h? Could you please have a look? Thanks very much.

@llvmbot
Copy link
Collaborator

llvmbot commented May 30, 2022

@llvm/issue-subscribers-opencl

@svenvh svenvh closed this as completed in a5cf17f May 30, 2022
@svenvh
Copy link
Member

svenvh commented May 30, 2022

I've pushed a fix for the immediate issue as a5cf17f . But I noticed some other discrepancies. @AnastasiaStulova I wonder if we should merge the OpenCL C 2.0 and 3.0 feature macro blocks in clang/lib/Headers/opencl-c-base.h ?

@AnastasiaStulova
Copy link
Contributor

I've pushed a fix for the immediate issue as a5cf17f . But I noticed some other discrepancies. @AnastasiaStulova I wonder if we should merge the OpenCL C 2.0 and 3.0 feature macro blocks in clang/lib/Headers/opencl-c-base.h ?

Makes sense. Otherwise we will keep fixing the same issues in multiple places!

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
Projects
None yet
Development

No branches or pull requests

5 participants