caio/clang-p2996
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

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

Browse Source 
Allow passing all zeros to reqd_work_group_size.

Test plan: ninja check-all
This commit is contained in:
Alexander Shaposhnikov
2025-03-16 16:21:46 -07:00
committed by GitHub
parent 81ba006296
commit 2dc123b33d
3 changed files with 18 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files

14
clang/lib/Sema/SemaDeclAttr.cpp
View File

@@ -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;
}
}
}

6
clang/test/CodeGenCUDASPIRV/spirv-attrs.cu
View File

@@ -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() {}
@@ -15,6 +18,8 @@ __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]+]]
@@ -22,6 +27,7 @@ __global__ void intel_reqd_sub_group_size_64() {}
// 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}

2
clang/test/SemaOpenCL/invalid-kernel-attrs.cl
View File

@@ -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(){}