From ee6bbb0004a35156b8f4bd0532ca711f2ee91799 Mon Sep 17 00:00:00 2001
From: sdashmiz <shadi.dashmiz@amd.com>
Date: Thu, 13 Jun 2024 16:04:17 -0400
Subject: [PATCH] FEAT-60269 - Test for large LDS

Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: I8f476f1b44324fffcc25423209a632caeb48cb99
---
 catch/include/hip_test_common.hh      | 20 ++++++++++++++++++++
 catch/unit/kernel/hipDynamicShared.cc | 12 +++++++++++-
 2 files changed, 31 insertions(+), 1 deletion(-)

diff --git a/catch/include/hip_test_common.hh b/catch/include/hip_test_common.hh
index 06be166269..0d4fb0c504 100644
--- a/catch/include/hip_test_common.hh
+++ b/catch/include/hip_test_common.hh
@@ -206,6 +206,26 @@ static inline bool IsGfx11() {
 #endif
 }
 
+static inline bool IsMi350() {
+#if HT_NVIDIA
+  return false;
+#elif HT_AMD
+  int device = -1;
+  hipDeviceProp_t props{};
+  HIP_CHECK(hipGetDevice(&device));
+  HIP_CHECK(hipGetDeviceProperties(&props, device));
+  // Get GCN Arch Name and compare to check if it is gfx11
+  std::string arch = std::string(props.gcnArchName);
+  auto pos = arch.find("gfx950");
+  if (pos != std::string::npos)
+    return true;
+  else
+    return false;
+#else
+  std::cout << "Have to be either Nvidia or AMD platform, asserting" << std::endl;
+  assert(false);
+#endif
+}
 
 // Utility Functions
 namespace HipTest {
diff --git a/catch/unit/kernel/hipDynamicShared.cc b/catch/unit/kernel/hipDynamicShared.cc
index 02a4ed9f19..8f9324ba2a 100644
--- a/catch/unit/kernel/hipDynamicShared.cc
+++ b/catch/unit/kernel/hipDynamicShared.cc
@@ -86,7 +86,7 @@ void testExternShared(size_t N, unsigned groupElements) {
   HIP_CHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
 
   // calculate the amount of dynamic shared memory required
-  size_t groupMemBytes = groupElements * sizeof(double);
+  size_t groupMemBytes = groupElements * sizeof(T);
 
   // launch kernel with dynamic shared memory
   hipLaunchKernelGGL(HIP_KERNEL_NAME(testExternSharedKernel<T>), dim3(blocks),
@@ -161,6 +161,16 @@ TEST_CASE("Unit_hipDynamicShared") {
     testExternShared<double>(65536, 32);
     testExternShared<double>(65536, 64);
   }
+
+  SECTION("test case with float for max LDS size") {
+    if(IsMi350())
+    {
+      testExternShared<float>(1024, 160*1024/sizeof(float));
+    }else
+    {
+      testExternShared<float>(1024, 64*1024/sizeof(float));
+    }
+  }
 }
 
 /**