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<SYCLReqdWorkGroupSizeAttr>()) {
-      auto GetDim = [](std::optional<llvm::APSInt> V) -> unsigned {
-        return V ? (unsigned)V->getZExtValue() : 1;
+      auto GetDim = [](std::optional<llvm::APSInt> 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);
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<class kernel_overflow>(f);
+  });
+  return 0;
+}
+
+// CHECK: define {{.*}} void @{{.*}}kernel_overflow() #[[ATTR:[0-9]+]]
+// CHECK: attributes #[[ATTR]] = { {{.*}}"amdgpu-flat-work-group-size"="4294967296,4294967296"{{.*}} }