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: 5b739b0373]

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),

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);