From 7ef0e834995db187bb3d93e533c41724873acfef Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Tue, 24 Nov 2020 02:34:34 +0000
Subject: [PATCH] Handle overflow

Fixes SWDEV-261712

Change-Id: I730fcdca8c2640fde9cfdfe2f4b4cf4894807fed


[ROCm/clr commit: 584771e7a25c7f140ddbace89408a9edcfdd6659]
---
 projects/clr/hipamd/rocclr/hip_module.cpp   | 18 +++++++++---------
 projects/clr/hipamd/rocclr/hip_platform.cpp |  6 +++---
 2 files changed, 12 insertions(+), 12 deletions(-)

diff --git a/projects/clr/hipamd/rocclr/hip_module.cpp b/projects/clr/hipamd/rocclr/hip_module.cpp
index 1480a3dc8c..06045e47bd 100755
--- a/projects/clr/hipamd/rocclr/hip_module.cpp
+++ b/projects/clr/hipamd/rocclr/hip_module.cpp
@@ -350,9 +350,9 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
                blockDimX, blockDimY, blockDimZ,
                sharedMemBytes, hStream,
                kernelParams, extra);
-  size_t globalWorkSizeX = gridDimX * blockDimX;
-  size_t globalWorkSizeY = gridDimY * blockDimY;
-  size_t globalWorkSizeZ = gridDimZ * blockDimZ;
+  size_t globalWorkSizeX = static_cast<size_t>(gridDimX) * blockDimX;
+  size_t globalWorkSizeY = static_cast<size_t>(gridDimY) * blockDimY;
+  size_t globalWorkSizeZ = static_cast<size_t>(gridDimZ) * blockDimZ;
   if (globalWorkSizeX > std::numeric_limits<uint32_t>::max() ||
       globalWorkSizeY > std::numeric_limits<uint32_t>::max() ||
       globalWorkSizeZ > std::numeric_limits<uint32_t>::max()) {
@@ -452,9 +452,9 @@ hipError_t hipLaunchCooperativeKernel(const void* f,
   int deviceId = ihipGetDevice();
   hipFunction_t func = nullptr;
   HIP_RETURN_ONFAIL(PlatformState::instance().getStatFunc(&func, f, deviceId));
-  size_t globalWorkSizeX = gridDim.x * blockDim.x;
-  size_t globalWorkSizeY = gridDim.y * blockDim.y;
-  size_t globalWorkSizeZ = gridDim.z * blockDim.z;
+  size_t globalWorkSizeX = static_cast<size_t>(gridDim.x) * blockDim.x;
+  size_t globalWorkSizeY = static_cast<size_t>(gridDim.y) * blockDim.y;
+  size_t globalWorkSizeZ = static_cast<size_t>(gridDim.z) * blockDim.z;
   if (globalWorkSizeX > std::numeric_limits<uint32_t>::max() ||
       globalWorkSizeY > std::numeric_limits<uint32_t>::max() ||
       globalWorkSizeZ > std::numeric_limits<uint32_t>::max()) {
@@ -539,9 +539,9 @@ hipError_t ihipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsL
       result = hipErrorInvalidDeviceFunction;
       HIP_RETURN(result);
     }
-    size_t globalWorkSizeX = launch.gridDim.x * launch.blockDim.x;
-    size_t globalWorkSizeY = launch.gridDim.y * launch.blockDim.y;
-    size_t globalWorkSizeZ = launch.gridDim.z * launch.blockDim.z;
+    size_t globalWorkSizeX = static_cast<size_t>(launch.gridDim.x) * launch.blockDim.x;
+    size_t globalWorkSizeY = static_cast<size_t>(launch.gridDim.y) * launch.blockDim.y;
+    size_t globalWorkSizeZ = static_cast<size_t>(launch.gridDim.z) * launch.blockDim.z;
     if (globalWorkSizeX > std::numeric_limits<uint32_t>::max() ||
         globalWorkSizeY > std::numeric_limits<uint32_t>::max() ||
         globalWorkSizeZ > std::numeric_limits<uint32_t>::max()) {
diff --git a/projects/clr/hipamd/rocclr/hip_platform.cpp b/projects/clr/hipamd/rocclr/hip_platform.cpp
index af4d51ef9a..4a58877ebe 100755
--- a/projects/clr/hipamd/rocclr/hip_platform.cpp
+++ b/projects/clr/hipamd/rocclr/hip_platform.cpp
@@ -607,9 +607,9 @@ hipError_t ihipLaunchKernel(const void* hostFunction,
   if ((hip_error != hipSuccess) || (func == nullptr)) {
     HIP_RETURN(hipErrorInvalidDeviceFunction);
   }
-  size_t globalWorkSizeX = gridDim.x * blockDim.x;
-  size_t globalWorkSizeY = gridDim.y * blockDim.y;
-  size_t globalWorkSizeZ = gridDim.z * blockDim.z;
+  size_t globalWorkSizeX = static_cast<size_t>(gridDim.x) * blockDim.x;
+  size_t globalWorkSizeY = static_cast<size_t>(gridDim.y) * blockDim.y;
+  size_t globalWorkSizeZ = static_cast<size_t>(gridDim.z) * blockDim.z;
   if (globalWorkSizeX > std::numeric_limits<uint32_t>::max() ||
       globalWorkSizeY > std::numeric_limits<uint32_t>::max() ||
       globalWorkSizeZ > std::numeric_limits<uint32_t>::max()) {