FEAT-60269 - Test for large LDS

Signed-off-by: sdashmiz <shadi.dashmiz@amd.com>
Change-Id: I8f476f1b44324fffcc25423209a632caeb48cb99

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 {

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));
+    }
+  }
 }
 
 /**