From 7ee9099c3efb952d4b86cc2f0b0016df3103341c Mon Sep 17 00:00:00 2001
From: Tao Sang <tao.sang@amd.com>
Date: Thu, 29 Jul 2021 16:55:10 -0400
Subject: [PATCH] SWDEV-294586 - Fix hipDynamicShared failure on NV

1. In kernel/hipDynamicShared
 Fix shared memory size and type mismatch in host and kernel.

2. In kernel/hipDynamicShared2
Cuda kernels relying on shared memory allocations over 48 KB require
to explicitly set size using hipFuncSetAttribute().

Change-Id: I4248b6cebd3dc156f9d5d427e1897da22fb964ed


[ROCm/hip commit: 5b739b0373993de78118135fd4b9618c22fd092c]
---
 .../hip/tests/src/kernel/hipDynamicShared.cpp  |  7 ++++---
 .../hip/tests/src/kernel/hipDynamicShared2.cpp | 18 +++++++++++++++---
 2 files changed, 19 insertions(+), 6 deletions(-)

diff --git a/projects/hip/tests/src/kernel/hipDynamicShared.cpp b/projects/hip/tests/src/kernel/hipDynamicShared.cpp
index 9532e0eaba..4552761d2f 100644
--- a/projects/hip/tests/src/kernel/hipDynamicShared.cpp
+++ b/projects/hip/tests/src/kernel/hipDynamicShared.cpp
@@ -22,7 +22,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * TEST: %t EXCLUDE_HIP_PLATFORM nvidia
+ * TEST: %t
  * HIT_END
  */
 
@@ -49,7 +49,8 @@ template <typename T>
 __global__ void testExternSharedKernel(const T* A_d, const T* B_d, T* C_d,
                                        size_t numElements, size_t groupElements) {
     // declare dynamic shared memory
-    extern __shared__ double sdata[];
+    extern __shared__ double sdata0[];
+    T* sdata = reinterpret_cast<T *>(sdata0);
 
     size_t gid = (blockIdx.x * blockDim.x + threadIdx.x);
     size_t tid = threadIdx.x;
@@ -90,7 +91,7 @@ void testExternShared(size_t N, unsigned groupElements) {
     HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice));
 
     // calculate the amount of dynamic shared memory required
-    size_t groupMemBytes = groupElements * sizeof(T);
+    size_t groupMemBytes = groupElements * sizeof(double);
 
     // launch kernel with dynamic shared memory
     hipLaunchKernelGGL(HIP_KERNEL_NAME(testExternSharedKernel<T>), dim3(blocks), dim3(threadsPerBlock),
diff --git a/projects/hip/tests/src/kernel/hipDynamicShared2.cpp b/projects/hip/tests/src/kernel/hipDynamicShared2.cpp
index 6c22324d8c..bca1fff722 100644
--- a/projects/hip/tests/src/kernel/hipDynamicShared2.cpp
+++ b/projects/hip/tests/src/kernel/hipDynamicShared2.cpp
@@ -22,15 +22,15 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../test_common.cpp
- * TEST: %t EXCLUDE_HIP_PLATFORM nvidia
+ * TEST: %t
  * HIT_END
  */
 
 #include "hip/hip_runtime.h"
 #include "test_common.h"
 
-#define LEN 16 * 1024
-#define SIZE LEN * 4
+#define LEN  (16 * 1024)
+#define SIZE (LEN * sizeof(float))
 
 __global__ void vectorAdd(float* Ad, float* Bd) {
     extern __shared__ float sBd[];
@@ -53,7 +53,19 @@ int main() {
     hipMalloc(&Bd, SIZE);
     hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
     hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice);
+
+    hipError_t ret = hipFuncSetAttribute(
+        reinterpret_cast<const void*>(&vectorAdd),
+        hipFuncAttributeMaxDynamicSharedMemorySize, SIZE);
+
+    if (ret != hipSuccess) {
+        printf("Failed requesting enough shared memory size(%zu), error: '%s'(%d), ignored!\n",
+             SIZE, hipGetErrorString(ret), ret);
+        passed();
+    }
+
     hipLaunchKernelGGL(vectorAdd, dim3(1, 1, 1), dim3(64, 1, 1), SIZE, 0, Ad, Bd);
+    HIPCHECK(hipGetLastError());
     hipMemcpy(B, Bd, SIZE, hipMemcpyDeviceToHost);
     for (int i = 0; i < LEN; i++) {
         assert(B[i] > 1.0f && B[i] < 3.0f);