Skip to content

[clang][opencl] Allow passing all zeros to reqd_work_group_size #131543

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.

Sign up for GitHub

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

Merged
merged 1 commit into from
Mar 16, 2025
Merged

[clang][opencl] Allow passing all zeros to reqd_work_group_size #131543

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
14 changes: 10 additions & 4 deletions clang/lib/Sema/SemaDeclAttr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2923,10 +2923,16 @@ static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
if (!S.checkUInt32Argument(AL, E, WGSize[i], i,
/*StrictlyUnsigned=*/true))
return;
if (WGSize[i] == 0) {
S.Diag(AL.getLoc(), diag::err_attribute_argument_is_zero)
<< AL << E->getSourceRange();
return;
}

if (!llvm::all_of(WGSize, [](uint32_t Size) { return Size == 0; })) {
for (unsigned i = 0; i < 3; ++i) {
const Expr *E = AL.getArgAsExpr(i);
if (WGSize[i] == 0) {
S.Diag(AL.getLoc(), diag::err_attribute_argument_is_zero)
<< AL << E->getSourceRange();
return;
}
}
}

Expand Down
6 changes: 6 additions & 0 deletions clang/test/CodeGenCUDASPIRV/spirv-attrs.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,6 +3,9 @@

#define __global__ __attribute__((global))

__attribute__((reqd_work_group_size(0, 0, 0)))
__global__ void reqd_work_group_size_0_0_0() {}

__attribute__((reqd_work_group_size(128, 1, 1)))
__global__ void reqd_work_group_size_128_1_1() {}

Expand All @@ -15,13 +18,16 @@ __global__ void vec_type_hint_int() {}
__attribute__((intel_reqd_sub_group_size(64)))
__global__ void intel_reqd_sub_group_size_64() {}


// CHECK: define spir_kernel void @_Z26reqd_work_group_size_0_0_0v() #[[ATTR:[0-9]+]] !reqd_work_group_size ![[WG_SIZE_ZEROS:[0-9]+]]
// CHECK: define spir_kernel void @_Z28reqd_work_group_size_128_1_1v() #[[ATTR:[0-9]+]] !reqd_work_group_size ![[WG_SIZE:[0-9]+]]
// CHECK: define spir_kernel void @_Z26work_group_size_hint_2_2_2v() #[[ATTR]] !work_group_size_hint ![[WG_HINT:[0-9]+]]
// CHECK: define spir_kernel void @_Z17vec_type_hint_intv() #[[ATTR]] !vec_type_hint ![[VEC_HINT:[0-9]+]]
// CHECK: define spir_kernel void @_Z28intel_reqd_sub_group_size_64v() #[[ATTR]] !intel_reqd_sub_group_size ![[SUB_GRP:[0-9]+]]

// CHECK: attributes #[[ATTR]] = { {{.*}} }

// CHECK: ![[WG_SIZE_ZEROS]] = !{i32 0, i32 0, i32 0}
// CHECK: ![[WG_SIZE]] = !{i32 128, i32 1, i32 1}
// CHECK: ![[WG_HINT]] = !{i32 2, i32 2, i32 2}
// CHECK: ![[VEC_HINT]] = !{i32 undef, i32 1}
Expand Down
2 changes: 2 additions & 0 deletions clang/test/SemaOpenCL/invalid-kernel-attrs.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,3 +44,5 @@ __kernel __attribute__((reqd_work_group_size(8,16,-32))) void neg2(){} // expect
// 4294967294 is a negative integer if treated as signed.
// Should compile successfully, since we expect an unsigned.
__kernel __attribute__((reqd_work_group_size(8,16,4294967294))) void ok1(){}

__kernel __attribute__((reqd_work_group_size(0,0,0))) void ok_zeros(){}
Loading