Adding lit tests

hipamd/CMakeLists.txt

@@ -1,17 +1,21 @@
 if(${HIPIFY_STANDLONE})
   cmake_minimum_required(VERSION 2.8.8)
-
+  
   project(hipify)
 
   find_package(LLVM REQUIRED PATHS ${LLVM_DIR} NO_DEFAULT_PATH)
 
+  list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR})
+  include(AddLLVM)
+
   message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}")
   message(STATUS "Using LLVMConfig.cmake in: ${LLVM_DIR}")
 
   include_directories(${LLVM_INCLUDE_DIRS})
   link_directories(${LLVM_LIBRARY_DIRS})
   add_definitions(${LLVM_DEFINITIONS})
-  add_executable( hipify src/Cuda2Hip.cpp )
+  add_llvm_executable( hipify src/Cuda2Hip.cpp )
+  
 else()
   set(LLVM_LINK_COMPONENTS
     Option
@@ -48,6 +52,7 @@ if(${HIPIFY_STANDLONE})
   set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}")
   set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS}")
   set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wall -std=c++11 -pthread -fno-rtti -fvisibility-inlines-hidden")
+  add_subdirectory(test)
 else()
   install(TARGETS hipify
           RUNTIME DESTINATION bin

hipamd/README.md

@@ -9,7 +9,7 @@ git clone http://llvm.org/git/llvm.git llvm
 git clone http://llvm.org/git/clang.git llvm/tools/clang
 
 mkdir llvm_build && cd llvm_build
-cmake -DCMAKE_INSTALL_PREFIX="LLVM_INSTALL_PATH" -DLLVM_TARGETS_TO_BUILD="X86;NVPTX;AMDGPU" ../llvm
+cmake -DCMAKE_INSTALL_PREFIX="LLVM_INSTALL_PATH" -DLLVM_INSTALL_UTILS=ON -DLLVM_TARGETS_TO_BUILD="X86;NVPTX;AMDGPU" ../llvm
 make && make install
 
 git clone https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP-hipify.git path_to_hipify_src
@@ -34,3 +34,14 @@ cmake -DLLVM_TARGETS_TO_BUILD="X86;NVPTX;AMDGPU" ../llvm
 make -C tools/clang/tools/extra/hipify
 ```
 
+# How to run tests (for *standalone* tool only)
+- install Python and add python-setuptools
+- install lit python script
+- make sure that FileCheck util is installed to **LLVM_INSTALL_PATH/bin/FileCheck**
+- run tests from path_to_hipify_src/build
+```
+sudo apt-get install python python-setuptools
+sudo easy_install lit
+make -C path_to_hipify_src/build test
+```
+

hipamd/test/CMakeLists.txt

@@ -0,0 +1,27 @@
+set(Python_ADDITIONAL_VERSIONS 2.7)
+include(FindPythonInterp)
+if( NOT PYTHONINTERP_FOUND )
+  message(FATAL_ERROR
+          "Unable to find Python interpreter, required for builds and testing\n\n"
+          "Please install Python or specify the PYTHON_EXECUTABLE CMake variable.")
+endif()
+
+set(BINARY_DIR       ${CMAKE_CURRENT_BINARY_DIR}  )
+
+configure_file(
+  ${CMAKE_CURRENT_SOURCE_DIR}/lit.site.cfg.in
+  ${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg
+  @ONLY)
+
+add_lit_testsuite(check-hipify "Running HIPify regression tests"
+  ${CMAKE_CURRENT_SOURCE_DIR}
+  PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/lit.site.cfg
+  DEPENDS hipify
+  )
+
+add_custom_target(check)
+add_dependencies(check check-hipify)
+add_custom_target(test)
+add_dependencies(test check-hipify)
+set_target_properties(check PROPERTIES FOLDER "Tests")
+

hipamd/test/axpy.cu

@@ -0,0 +1,43 @@
+// RUN: hipify "%s" 2>&1 | FileCheck %s
+
+#include <helper_cuda.h> // for checkCudaErrors
+
+#include <iostream>
+
+__global__ void axpy(float a, float* x, float* y) {
+  // CHECK: hipThreadIdx_x
+  y[threadIdx.x] = a * x[threadIdx.x];
+}
+
+int main(int argc, char* argv[]) {
+  const int kDataLen = 4;
+
+  float a = 2.0f;
+  float host_x[kDataLen] = {1.0f, 2.0f, 3.0f, 4.0f};
+  float host_y[kDataLen];
+
+  // Copy input data to device.
+  float* device_x;
+  float* device_y;
+  checkCudaErrors(cudaMalloc(&device_x, kDataLen * sizeof(float)));
+  checkCudaErrors(cudaMalloc(&device_y, kDataLen * sizeof(float)));
+  checkCudaErrors(cudaMemcpy(device_x, host_x, kDataLen * sizeof(float),
+                             cudaMemcpyHostToDevice));
+
+  // Launch the kernel.
+  // CHECK: hipLaunchKernel(HIP_KERNEL_NAME
+  axpy<<<1, kDataLen>>>(a, device_x, device_y);
+
+  // Copy output data to host.
+  checkCudaErrors(cudaDeviceSynchronize());
+  checkCudaErrors(cudaMemcpy(host_y, device_y, kDataLen * sizeof(float),
+                             cudaMemcpyDeviceToHost));
+
+  // Print the results.
+  for (int i = 0; i < kDataLen; ++i) {
+    std::cout << "y[" << i << "] = " << host_y[i] << "\n";
+  }
+
+  checkCudaErrors(cudaDeviceReset());
+  return 0;
+}

hipamd/test/lit.cfg

@@ -0,0 +1,54 @@
+# -*- Python -*-
+import os
+import platform
+import re
+import subprocess
+
+import lit.formats
+import lit.util
+
+# Configuration file for the 'lit' test runner.
+
+# name: The name of this test suite.
+config.name = 'hipify'
+
+# suffixes: CUDA source is only supported
+config.suffixes = ['.cu']
+
+# testFormat: The test format to use to interpret tests.
+config.test_format = lit.formats.ShTest()
+
+# test_source_root: The root path where tests are located.
+config.test_source_root = os.path.dirname(__file__)
+
+# test_exec_root: The path where tests are located (default is the test suite root).
+#config.test_exec_root = config.test_source_root
+
+# target_triple: Used by ShTest and TclTest formats for XFAIL checks.
+config.target_triple = '(unused)'
+
+# available_features: Used by ShTest and TclTest formats for REQUIRES checks.
+config.available_features = []
+
+site_cfg = lit_config.params.get('site_config', None)
+lit_config.load_config(config, site_cfg)
+
+obj_root = getattr(config, 'obj_root', None)
+if obj_root is not None:
+    config.test_exec_root = obj_root
+
+if obj_root is not None:
+    llvm_tools_dir = getattr(config, 'llvm_tools_dir', None)
+    if not llvm_tools_dir:
+        lit_config.fatal('No LLVM tools dir set!')
+    path = os.path.pathsep.join((llvm_tools_dir, config.environment['PATH']))
+    config.environment['PATH'] = path
+
+tool_name = "FileCheck"
+tool_path = lit.util.which(tool_name, llvm_tools_dir)
+if not tool_path:
+    # Warn, but still provide a substitution.
+    lit_config.note('Did not find ' + tool_name + ' in ' + llvm_tools_dir)
+    tool_path = llvm_tools_dir + '/' + tool_name
+config.substitutions.append((tool_name, tool_path))
+

hipamd/test/lit.site.cfg.in

@@ -0,0 +1,15 @@
+import sys
+
+config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@"
+config.obj_root = "@CMAKE_CURRENT_BINARY_DIR@"
+
+# Support substitution of the tools and libs dirs with user parameters. This is
+# used when we can't determine the tool dir at configuration time.
+try:
+    config.llvm_tools_dir = config.llvm_tools_dir % lit_config.params
+    config.obj_root = config.obj_root % lit_config.params
+except KeyError:
+    e = sys.exc_info()[1]
+    key, = e.args
+    lit_config.fatal("unable to find %r parameter, use '--param=%s=VALUE'" % (key,key))
+