From 3ae3c39e44c3df4a5ac101bf35994f8fb74ff62f Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Fri, 14 Oct 2016 23:19:25 -0500
Subject: [PATCH] Refactor module API test.

- Add PASSED/FAIL indication.
- Set args using struct rather than void* array.

Change-Id: Ic924f88c49cc46979b12b7fef8650081e3b5f58c
---
 samples/0_Intro/module_api/runKernel.cpp | 55 +++++++++++++++---------
 1 file changed, 35 insertions(+), 20 deletions(-)

diff --git a/samples/0_Intro/module_api/runKernel.cpp b/samples/0_Intro/module_api/runKernel.cpp
index 5f16677fc2..90b081e09c 100644
--- a/samples/0_Intro/module_api/runKernel.cpp
+++ b/samples/0_Intro/module_api/runKernel.cpp
@@ -60,32 +60,37 @@ int main(){
 		uint32_t len = LEN;
 		uint32_t one = 1;
 
-    std::vector<void*>argBuffer(5);
-    uint32_t *ptr32_t = (uint32_t*)&argBuffer[0];
-    memcpy(ptr32_t + 0, &one, sizeof(uint32_t));
-    memcpy(ptr32_t + 1, &one, sizeof(uint32_t));
-    memcpy(ptr32_t + 2, &one, sizeof(uint32_t));
-    memcpy(ptr32_t + 3, &len, sizeof(uint32_t));
-    memcpy(ptr32_t + 4, &one, sizeof(uint32_t));
-    memcpy(ptr32_t + 5, &one, sizeof(uint32_t));
-    memcpy(&argBuffer[3], &Ad, sizeof(void*));
-    memcpy(&argBuffer[4], &Bd, sizeof(void*));
+    struct {
+        uint32_t _hidden[6];
+        void * _Ad;
+        void * _Bd;
+    } args;
+
+    for (int i=0; i<6; i++) {
+        args._hidden[i] = 0;
+    }
+    args._Ad = Ad;
+    args._Bd = Bd;
+
 #endif
 
 #ifdef __HIP_PLATFORM_NVCC__
-	uint32_t one = 1;
-    std::vector<void*>argBuffer(3);
-    uint32_t *ptr32_t = (uint32_t*)&argBuffer[0];
-    memcpy(ptr32_t + 0, &one, sizeof(uint32_t));
-    memcpy(&argBuffer[1], &Ad, sizeof(void*));
-    memcpy(&argBuffer[2], &Bd, sizeof(void*));
+    struct {
+        uint32_t _hidden[1];
+        void * _Ad;
+        void * _Bd;
+    } args;
+
+    args._hidden[0] = 0;
+    args._Ad = Ad;
+    args._Bd = Bd;
 #endif
 
 
-    size_t size = argBuffer.size()*sizeof(void*);
+    size_t size = sizeof(args);
 
     void *config[] = {
-      HIP_LAUNCH_PARAM_BUFFER_POINTER, &argBuffer[0],
+      HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
       HIP_LAUNCH_PARAM_BUFFER_SIZE, &size,
       HIP_LAUNCH_PARAM_END
     };
@@ -93,10 +98,20 @@ int main(){
     hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, (void**)&config); 
 
     hipMemcpyDtoH(B, Bd, SIZE);
-    for(uint32_t i=LEN-4;i<LEN;i++){
-        std::cout<<A[i]<<" - "<<B[i]<<std::endl;
+    int mismatchCount = 0;
+    for(uint32_t i=0;i<LEN;i++){
+        if (A[i] != B[i]) {
+            mismatchCount++;
+            std::cout<<"error: mismatch " << A[i]<<" != "<<B[i]<<std::endl;
+        }
     }
 
+    if (mismatchCount == 0) {
+        std::cout << "PASSED!\n";
+    } else {
+        std::cout << "FAILED!\n";
+    };
+
     hipCtxDestroy(context);
     return 0;
 }