From 8155ee0797478ffce2a7a601d7b43abf2f88e1e0 Mon Sep 17 00:00:00 2001 From: Wenju He Date: Tue, 14 Jul 2026 03:22:58 +0200 Subject: [PATCH 1/2] [SYCL][AMDGPU] Use uint64_t type for SYCLReqdWorkGroupSize product type range::size() is product of its dimensions and its return type is size_t. SYCLReqdWorkGroupSize product probably should fit in size_t type. unsigned type may have overflow issue when size_t type size is 64bit. Use uint64_t to align with downstream code. --- clang/lib/CodeGen/Targets/AMDGPU.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp index 69469704d1228..65c9ef2135163 100644 --- a/clang/lib/CodeGen/Targets/AMDGPU.cpp +++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp @@ -352,10 +352,10 @@ void AMDGPUTargetCodeGenInfo::setFunctionDeclAttributes( M.handleAMDGPUFlatWorkGroupSizeAttr(F, FlatWGS, ReqdWGS); } else if (M.getLangOpts().SYCLIsDevice) { if (const auto *SYCLReqdWGS = FD->getAttr()) { - auto GetDim = [](std::optional V) -> unsigned { - return V ? (unsigned)V->getZExtValue() : 1; + auto GetDim = [](std::optional V) -> uint64_t { + return V ? V->getZExtValue() : 1; }; - unsigned Size = GetDim(SYCLReqdWGS->getXDimVal()) * + uint64_t Size = GetDim(SYCLReqdWGS->getXDimVal()) * GetDim(SYCLReqdWGS->getYDimVal()) * GetDim(SYCLReqdWGS->getZDimVal()); std::string AttrVal = llvm::utostr(Size) + "," + llvm::utostr(Size); From 308764752f9fbcb2123f32e93ed1b2cf2482ee00 Mon Sep 17 00:00:00 2001 From: Wenju He Date: Fri, 17 Jul 2026 05:05:41 +0200 Subject: [PATCH 2/2] add test Co-authored-by: Claude Sonnet 5 --- .../reqd-work-group-size-overflow.cpp | 27 +++++++++++++++++++ 1 file changed, 27 insertions(+) create mode 100644 clang/test/CodeGenSYCL/reqd-work-group-size-overflow.cpp diff --git a/clang/test/CodeGenSYCL/reqd-work-group-size-overflow.cpp b/clang/test/CodeGenSYCL/reqd-work-group-size-overflow.cpp new file mode 100644 index 0000000000000..ded39def494e3 --- /dev/null +++ b/clang/test/CodeGenSYCL/reqd-work-group-size-overflow.cpp @@ -0,0 +1,27 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple amdgcn-amd-amdhsa -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s + +// The product of the reqd_work_group_size dimensions below is 2^32, which +// overflows a 32-bit unsigned integer. Check that the resulting +// amdgpu-flat-work-group-size function attribute reflects the full 64-bit +// product instead of a wrapped-around 32-bit value. + +#include "Inputs/sycl.hpp" + +using namespace sycl; +queue q; + +class FunctorOverflow { +public: + [[sycl::reqd_work_group_size(65536, 65536)]] void operator()() const {} +}; + +int main() { + q.submit([&](handler &h) { + FunctorOverflow f; + h.single_task(f); + }); + return 0; +} + +// CHECK: define {{.*}} void @{{.*}}kernel_overflow() #[[ATTR:[0-9]+]] +// CHECK: attributes #[[ATTR]] = { {{.*}}"amdgpu-flat-work-group-size"="4294967296,4294967296"{{.*}} }