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), dim3(blocks), @@ -161,6 +161,16 @@ TEST_CASE("Unit_hipDynamicShared") { testExternShared(65536, 32); testExternShared(65536, 64); } + + SECTION("test case with float for max LDS size") { + if(IsMi350()) + { + testExternShared(1024, 160*1024/sizeof(float)); + }else + { + testExternShared(1024, 64*1024/sizeof(float)); + } + } } /**