From 3f3e84b29e8fa3f5f091a7e775d5bdabe5d331c3 Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <Evgeniy.Mankov@amd.com>
Date: Wed, 18 Jul 2018 18:33:44 +0300
Subject: [PATCH] [HIPIFY][test] Add static and dynamic shared memory tests

[ROCm/hip commit: 70bd255d713cb180140cdd9064db0d1ca67217b0]
---
 .../hipify-clang/dynamic_shared_memory.cu     | 44 +++++++++++++++++++
 .../hipify-clang/static_shared_memory.cu      | 43 ++++++++++++++++++
 2 files changed, 87 insertions(+)
 create mode 100644 projects/hip/tests/hipify-clang/dynamic_shared_memory.cu
 create mode 100644 projects/hip/tests/hipify-clang/static_shared_memory.cu

diff --git a/projects/hip/tests/hipify-clang/dynamic_shared_memory.cu b/projects/hip/tests/hipify-clang/dynamic_shared_memory.cu
new file mode 100644
index 0000000000..d35718d714
--- /dev/null
+++ b/projects/hip/tests/hipify-clang/dynamic_shared_memory.cu
@@ -0,0 +1,44 @@
+// RUN: %run_test hipify "%s" "%t" %cuda_args
+
+// Taken from Jonathan Hui blog https://jhui.github.io/2017/03/06/CUDA
+
+#include <stdio.h>
+// CHECK: #include <hip/hip_runtime.h>
+#include <cuda.h>
+
+__global__ void dynamicReverse(int *d, int n)
+{
+  // Dynamic shared memory
+  // CHECK: HIP_DYNAMIC_SHARED(int, s);
+  extern __shared__ int s[];
+  int t = threadIdx.x;
+  int tr = n-t-1;
+  s[t] = d[t];
+  __syncthreads();
+  d[t] = s[tr];
+}
+
+int main(void)
+{
+  const int n = 64;
+  int a[n], r[n], d[n];
+
+  for (int i = 0; i < n; i++) {
+    a[i] = i;
+    r[i] = n-i-1;
+    d[i] = 0;
+  }
+
+  int *d_d;
+  // CHECK: hipMalloc(&d_d, n * sizeof(int));
+  cudaMalloc(&d_d, n * sizeof(int));
+  // run version with dynamic shared memory
+  // CHECK: hipMemcpy(d_d, a, n*sizeof(int), hipMemcpyHostToDevice);
+  cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);
+  // CHECK: hipLaunchKernelGGL(dynamicReverse, dim3(1), dim3(n), n*sizeof(int), 0, d_d, n);
+  dynamicReverse<<<1,n,n*sizeof(int)>>>(d_d, n);
+  // CHECK: hipMemcpy(d, d_d, n*sizeof(int), hipMemcpyDeviceToHost);
+  cudaMemcpy(d, d_d, n*sizeof(int), cudaMemcpyDeviceToHost);
+  for (int i = 0; i < n; i++)
+    if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)n", i, i, d[i], r[i]);
+}
diff --git a/projects/hip/tests/hipify-clang/static_shared_memory.cu b/projects/hip/tests/hipify-clang/static_shared_memory.cu
new file mode 100644
index 0000000000..ea6b50488f
--- /dev/null
+++ b/projects/hip/tests/hipify-clang/static_shared_memory.cu
@@ -0,0 +1,43 @@
+// RUN: %run_test hipify "%s" "%t" %cuda_args
+
+// Taken from Jonathan Hui blog https://jhui.github.io/2017/03/06/CUDA
+
+#include <stdio.h>
+// CHECK: #include <hip/hip_runtime.h>
+#include <cuda.h>
+
+__global__ void staticReverse(int *d, int n)
+{
+  __shared__ int s[64];
+  int t = threadIdx.x;
+  int tr = n-t-1;
+  s[t] = d[t];
+  // Will not conttinue until all threads completed.
+  __syncthreads();
+  d[t] = s[tr];
+}
+
+int main(void)
+{
+  const int n = 64;
+  int a[n], r[n], d[n];
+
+  for (int i = 0; i < n; i++) {
+    a[i] = i;
+    r[i] = n-i-1;
+    d[i] = 0;
+  }
+
+  int *d_d;
+  // CHECK: hipMalloc(&d_d, n * sizeof(int));
+  cudaMalloc(&d_d, n * sizeof(int));
+  // run version with static shared memory
+  // CHECK: hipMemcpy(d_d, a, n*sizeof(int), hipMemcpyHostToDevice);
+  cudaMemcpy(d_d, a, n*sizeof(int), cudaMemcpyHostToDevice);
+  // CHECK: hipLaunchKernelGGL(staticReverse, dim3(1), dim3(n), 0, 0, d_d, n);
+  staticReverse<<<1,n>>>(d_d, n);
+  // CHECK: hipMemcpy(d, d_d, n*sizeof(int), hipMemcpyDeviceToHost);
+  cudaMemcpy(d, d_d, n*sizeof(int), cudaMemcpyDeviceToHost);
+  for (int i = 0; i < n; i++)
+    if (d[i] != r[i]) printf("Error: d[%d]!=r[%d] (%d, %d)n", i, i, d[i], r[i]);
+}