From e6a4e455e55f84bd75e0c9d7fd2c16d53dfd0aed Mon Sep 17 00:00:00 2001
From: Tao Sang <tao.sang@amd.com>
Date: Mon, 2 Nov 2020 09:39:56 -0500
Subject: [PATCH] Make hipMemset*Async* tests pass in static lib

3 hipMemset*Async* direct tests will crash in static lib because
some hsa objects are not initialized when hip apis call them.
The fix will make the hip apis called in main() so that hsa objects
have been initialized at that stage.

Change-Id: I41eb29b8c2588acac0cd3ae6d96b14fb1abd235a


[ROCm/clr commit: 272f89a96611ba92e9b8e1c0b4455b338db8462d]
---
 .../memory/hipMemset2DAsyncMultiThreadAndKernel.cpp        | 7 +++++--
 .../src/runtimeApi/memory/hipMemsetAsyncAndKernel.cpp      | 7 +++++--
 .../src/runtimeApi/memory/hipMemsetAsyncMultiThread.cpp    | 7 ++++---
 3 files changed, 14 insertions(+), 7 deletions(-)

diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset2DAsyncMultiThreadAndKernel.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset2DAsyncMultiThreadAndKernel.cpp
index 4e3dfad062..a9097e0bfe 100644
--- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset2DAsyncMultiThreadAndKernel.cpp
+++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemset2DAsyncMultiThreadAndKernel.cpp
@@ -24,7 +24,7 @@
 //
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_LIB_TYPE static
+ * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
  * TEST: %t
  * HIT_END
  */
@@ -35,7 +35,7 @@
 #define NUM_H 256
 #define NUM_W 256
 
-unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
+unsigned blocks = 0;
 hipStream_t stream;
 
 bool testResult = true;
@@ -150,6 +150,9 @@ bool testhipMemset2DAsyncMultiThread() {
 int main() {
   bool testResult = true;
 
+  blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
+  printf("blocks: %u\n", blocks);
+
   testResult &= testhipMemset2DAsyncWithKernel();
   if (testResult) {
     printf("Kernel and hipMemset2DAsync executed in correct order!\n");
diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetAsyncAndKernel.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetAsyncAndKernel.cpp
index a24c9459a6..e3c2593c23 100644
--- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetAsyncAndKernel.cpp
+++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetAsyncAndKernel.cpp
@@ -23,7 +23,7 @@
  */
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_LIB_TYPE static
+ * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
  * TEST: %t
  * HIT_END
  */
@@ -32,7 +32,7 @@
 #define ITER 10
 #define N 1024 * 1024
 
-unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
+unsigned blocks = 0;
 
 template <typename T>
 __global__ void vector_square(T* B_d, T* C_d, size_t M) {
@@ -153,6 +153,9 @@ bool testhipMemsetD8AsyncWithKernel() {
 int main() {
   bool testResult = true;
   int numDevices = 0;
+  blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
+  printf("blocks: %u\n", blocks);
+
   HIPCHECK(hipGetDeviceCount(&numDevices));
   printf("total number of gpus in the system: %d\n", numDevices);
 
diff --git a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetAsyncMultiThread.cpp b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetAsyncMultiThread.cpp
index 490aef8c68..071554f661 100644
--- a/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetAsyncMultiThread.cpp
+++ b/projects/clr/hipamd/tests/src/runtimeApi/memory/hipMemsetAsyncMultiThread.cpp
@@ -22,7 +22,7 @@
  */
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_LIB_TYPE static
+ * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvidia
  * TEST: %t
  * HIT_END
  */
@@ -31,8 +31,6 @@
 #define NUM_THREADS 50
 #define ITER 50
 
-unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
-
 template <typename T>
 class MemSetTest {
  public:
@@ -219,6 +217,9 @@ int main() {
   printf("Queueing up hipMemSetAsync jobs on multiple threads"
          "and checking results\n");
 
+  unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N);
+  printf("blocks: %u\n", blocks);
+
   testResult &= testhipMemsetAsyncWithMultiThread();
   if (!(testResult)) {
     printf("Thread execution did not complete for hipMemsetAsync\n");