From 71b93692563e5b06c4e797044f6514dc8dbab422 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Thu, 16 Jun 2016 08:41:32 -0500
Subject: [PATCH] Add first steps for CMake test hierarchy + initial
 launch_bounds.

---
 hipamd/tests/src/CMakeLists.txt               | 13 +++-
 hipamd/tests/src/launch_bounds/CMakeLists.txt |  7 ++
 .../src/launch_bounds/hip_launch_bounds.cpp   | 73 +++++++++++++++++++
 3 files changed, 89 insertions(+), 4 deletions(-)
 create mode 100644 hipamd/tests/src/launch_bounds/CMakeLists.txt
 create mode 100644 hipamd/tests/src/launch_bounds/hip_launch_bounds.cpp

diff --git a/hipamd/tests/src/CMakeLists.txt b/hipamd/tests/src/CMakeLists.txt
index d758d3c54f..f1e37d242b 100644
--- a/hipamd/tests/src/CMakeLists.txt
+++ b/hipamd/tests/src/CMakeLists.txt
@@ -2,7 +2,9 @@ cmake_minimum_required (VERSION 2.6)
 project (HIP_Unit_Tests)
 include(CTest)
 
-include_directories( ${PROJECT_SOURCE_DIR}/include )
+
+#include_directories( ${PROJECT_SOURCE_DIR}/include )
+set (HIPTEST_SOURCE_DIR ${PROJECT_SOURCE_DIR} )
 
 # The version number.
 set (HIP_Unit_Test_VERSION_MAJOR 1)
@@ -22,7 +24,7 @@ endif()
 
 set(HIP_PATH $ENV{HIP_PATH})
 if (NOT DEFINED HIP_PATH)
-	set (HIP_PATH ../..)
+    get_filename_component (HIP_PATH ../.. ABSOLUTE)
 endif()
 
 execute_process(COMMAND ${HIP_PATH}/bin/hipconfig --platform OUTPUT_VARIABLE HIP_PLATFORM)
@@ -105,14 +107,14 @@ macro (make_hip_executable_libcpp exe cpp)
     endif()
 endmacro()
 
-macro (make_named_test exe testname )
+function (make_named_test exe testname )
   add_test (NAME ${testname}
       COMMAND ${PROJECT_BINARY_DIR}/${exe} ${ARGN}
             )
   set_tests_properties (${testname}
       PROPERTIES PASS_REGULAR_EXPRESSION "PASSED"
     )
-endmacro()
+endfunction()
 
 macro (make_test exe )
   string (REPLACE " " "" smush_args ${ARGN})
@@ -263,3 +265,6 @@ endif()
 make_hipify_test(specialFunc.cu )
 
 make_test(hipDynamicShared " ")
+
+# Add subdirs here:
+add_subdirectory(launch_bounds)
diff --git a/hipamd/tests/src/launch_bounds/CMakeLists.txt b/hipamd/tests/src/launch_bounds/CMakeLists.txt
new file mode 100644
index 0000000000..d9d96ee625
--- /dev/null
+++ b/hipamd/tests/src/launch_bounds/CMakeLists.txt
@@ -0,0 +1,7 @@
+cmake_minimum_required (VERSION 2.6)
+project ( ${PROJECT_SOURCE_DIR}/launch_bounds)
+
+include_directories( ${HIPTEST_SOURCE_DIR} )
+
+make_hip_executable (hip_launch_bounds hip_launch_bounds.cpp)
+make_test(hip_launch_bounds " ")
diff --git a/hipamd/tests/src/launch_bounds/hip_launch_bounds.cpp b/hipamd/tests/src/launch_bounds/hip_launch_bounds.cpp
new file mode 100644
index 0000000000..cba9af27c8
--- /dev/null
+++ b/hipamd/tests/src/launch_bounds/hip_launch_bounds.cpp
@@ -0,0 +1,73 @@
+/*
+Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+// Simple test for memset.
+// Also serves as a template for other tests.
+
+#include "hip_runtime.h"
+#include "test_common.h"
+
+
+__global__ 
+void
+myKern(hipLaunchParm lp, int *C, const int *A, int N)
+{
+    int tid = hipThreadIdx_x;
+
+    C[tid] = A[tid];
+};
+
+
+
+int main(int argc, char *argv[])
+{
+    HipTest::parseStandardArguments(argc, argv, true);
+
+    size_t Nbytes = N*sizeof(int);
+
+    int *A_d, *C_d, *A_h, *C_h;
+    HIPCHECK ( hipMalloc(&A_d, Nbytes) );
+    HIPCHECK ( hipMalloc(&C_d, Nbytes) );
+    A_h = (int*)malloc (Nbytes);
+
+    C_h = (int*)malloc (Nbytes);
+    for (int i=0; i<N; i++) {
+        A_h[i] = i*10;
+        C_h[i] = 0x0;
+    }
+
+    HIPCHECK ( hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice) );
+
+    hipLaunchKernel(myKern, dim3(N), dim3(256), 0, 0, A_d, C_d, N);
+
+    HIPCHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost) );
+
+    for (int i=0; i<N; i++) {
+        int goldVal = i * 10;
+        if (A_h[i] != goldVal) {
+            failed("mismatch at index:%d computed:%02x, gold:%02x\n", i, (int)A_h[i], (int)goldVal);
+
+        }
+    }
+
+    passed();
+
+};