From 0a918c8f96097f9417eea7d14c9a998e61aedb2b Mon Sep 17 00:00:00 2001
From: pghafari <pghafari@amd.com>
Date: Fri, 6 Sep 2024 19:04:21 -0400
Subject: [PATCH] SWDEV-479260,SWDEV-483599 - Check griddim Y,Z <= 65536

Gfx12 has 16 bits for grid dim Y/Z. Detect gfxIp and return error if dim y/z > 16 bits

Change-Id: I43dd14affc9e4073d0b1232e7523967f0180fa31
---
 hipamd/src/hip_platform.cpp | 7 +++++++
 1 file changed, 7 insertions(+)

diff --git a/hipamd/src/hip_platform.cpp b/hipamd/src/hip_platform.cpp
index 535d6e96f6..11b00ed735 100644
--- a/hipamd/src/hip_platform.cpp
+++ b/hipamd/src/hip_platform.cpp
@@ -619,6 +619,13 @@ hipError_t ihipLaunchKernel(const void* hostFunction, dim3 gridDim, dim3 blockDi
       return hipErrorInvalidDeviceFunction;
     }
   }
+
+  constexpr auto gridDimYZmax = static_cast<uint64_t>(std::numeric_limits<uint16_t>::max()) + 1;
+  const auto& isa = g_devices[deviceId]->devices()[0]->isa().versionMajor();
+  if (isa >= 12 && (gridDim.y > gridDimYZmax || gridDim.z > gridDimYZmax)) {
+    return hipErrorInvalidConfiguration;
+  }
+
   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;