From a0720fe79afa87063d33bb5163913944e3b083f4 Mon Sep 17 00:00:00 2001 From: dfukalov Date: Wed, 17 Feb 2016 19:05:18 +0300 Subject: [PATCH 1/2] Adding lit tests --- CMakeLists.txt | 9 ++++++-- README.md | 13 ++++++++++- test/CMakeLists.txt | 27 ++++++++++++++++++++++ test/axpy.cu | 43 +++++++++++++++++++++++++++++++++++ test/lit.cfg | 54 ++++++++++++++++++++++++++++++++++++++++++++ test/lit.site.cfg.in | 15 ++++++++++++ 6 files changed, 158 insertions(+), 3 deletions(-) create mode 100644 test/CMakeLists.txt create mode 100644 test/axpy.cu create mode 100644 test/lit.cfg create mode 100644 test/lit.site.cfg.in diff --git a/CMakeLists.txt b/CMakeLists.txt index 56845f1b6e..8915cc4714 100644 --- a/CMakeLists.txt +++ b/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 diff --git a/README.md b/README.md index 08b3a5eb5f..a3725394ad 100644 --- a/README.md +++ b/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 +``` + diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt new file mode 100644 index 0000000000..4f74091b0b --- /dev/null +++ b/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") + diff --git a/test/axpy.cu b/test/axpy.cu new file mode 100644 index 0000000000..29d5493763 --- /dev/null +++ b/test/axpy.cu @@ -0,0 +1,43 @@ +// RUN: hipify "%s" 2>&1 | FileCheck %s + +#include // for checkCudaErrors + +#include + +__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; +} diff --git a/test/lit.cfg b/test/lit.cfg new file mode 100644 index 0000000000..775cdbf0d4 --- /dev/null +++ b/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)) + diff --git a/test/lit.site.cfg.in b/test/lit.site.cfg.in new file mode 100644 index 0000000000..6fc35440d4 --- /dev/null +++ b/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)) + From f3dfd07fa57af14a6ea007767598ba6eeda91fb9 Mon Sep 17 00:00:00 2001 From: dfukalov Date: Thu, 18 Feb 2016 23:16:52 +0300 Subject: [PATCH 2/2] fix build bug with current clang/llvm --- src/Cuda2Hip.cpp | 76 ++++++++++++++++++++++++------------------------ 1 file changed, 38 insertions(+), 38 deletions(-) diff --git a/src/Cuda2Hip.cpp b/src/Cuda2Hip.cpp index fe3f360dce..f8cf938f32 100644 --- a/src/Cuda2Hip.cpp +++ b/src/Cuda2Hip.cpp @@ -229,22 +229,22 @@ namespace { DenseMap cuda2hipRename; }; - struct HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks { - HipifyPPCallbacks(Replacements * R) - : SeenEnd(false), _sm(nullptr), Replace(R) - { - } - - virtual bool handleBeginSource(CompilerInstance &CI, StringRef Filename) override - { - Preprocessor &PP = CI.getPreprocessor(); - SourceManager & SM = CI.getSourceManager(); - setSourceManager(&SM); - PP.addPPCallbacks(std::unique_ptr(this)); - PP.Retain(); - return true; - } - + struct HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks { + HipifyPPCallbacks(Replacements * R) + : SeenEnd(false), _sm(nullptr), Replace(R) + { + } + + virtual bool handleBeginSource(CompilerInstance &CI, StringRef Filename) override + { + Preprocessor &PP = CI.getPreprocessor(); + SourceManager & SM = CI.getSourceManager(); + setSourceManager(&SM); + PP.addPPCallbacks(std::unique_ptr(this)); + PP.Retain(); + return true; + } + virtual void InclusionDirective( SourceLocation hash_loc, const Token &include_token, @@ -254,7 +254,7 @@ namespace { const FileEntry *file, StringRef search_path, StringRef relative_path, - const Module *imported) override + const clang::Module *imported) override { if (_sm->isWrittenInMainFile(hash_loc)) { if (is_angled) { @@ -265,8 +265,8 @@ namespace { llvm::errs() << "\nWill be replaced with " << repName << "\n"; SourceLocation sl = filename_range.getBegin(); SourceLocation sle = filename_range.getEnd(); - const char* B = _sm->getCharacterData(sl); - const char* E = _sm->getCharacterData(sle); + const char* B = _sm->getCharacterData(sl); + const char* E = _sm->getCharacterData(sle); repName = "<" + repName + ">"; Replacement Rep(*_sm, sl, E - B, repName); Replace->insert(Rep); @@ -300,22 +300,22 @@ namespace { } } } - - void EndOfMainFile() override - { - - } - - bool SeenEnd; - void setSourceManager(SourceManager * sm) { _sm = sm; } - - private: - - SourceManager * _sm; - Replacements * Replace; - struct hipName N; - }; - + + void EndOfMainFile() override + { + + } + + bool SeenEnd; + void setSourceManager(SourceManager * sm) { _sm = sm; } + + private: + + SourceManager * _sm; + Replacements * Replace; + struct hipName N; + }; + class Cuda2HipCallback : public MatchFinder::MatchCallback { public: Cuda2HipCallback(Replacements *Replace) : Replace(Replace) {} @@ -452,8 +452,8 @@ class Cuda2HipCallback : public MatchFinder::MatchCallback { if (const ParmVarDecl * cudaParamDecl = Result.Nodes.getNodeAs("cudaParamDecl")) { - QualType QT = cudaParamDecl->getOriginalType(); - std::string name = QT.getAsString(); + QualType QT = cudaParamDecl->getOriginalType(); + std::string name = QT.getAsString(); std::string repName = N.cuda2hipRename[name]; SourceLocation sl = cudaParamDecl->getLocStart(); Replacement Rep(*SM, SM->isMacroArgExpansion(sl) ? @@ -506,7 +506,7 @@ int main(int argc, const char **argv) { Finder.addMatcher(declRefExpr(isExpansionInMainFile(), to(enumConstantDecl(matchesName("cuda.*")))).bind("cudaEnumConstantRef"), &Callback); Finder.addMatcher(varDecl(isExpansionInMainFile(), hasType(enumDecl(matchesName("cuda.*")))).bind("cudaEnumConstantDecl"), &Callback); Finder.addMatcher(varDecl(isExpansionInMainFile(), hasType(cxxRecordDecl(matchesName("cuda.*")))).bind("cudaStructVar"), &Callback); - Finder.addMatcher(parmVarDecl(isExpansionInMainFile(), hasType(namedDecl(matchesName("cuda.*")))).bind("cudaParamDecl"), &Callback); + Finder.addMatcher(parmVarDecl(isExpansionInMainFile(), hasType(namedDecl(matchesName("cuda.*")))).bind("cudaParamDecl"), &Callback); auto action = newFrontendActionFactory(&Finder, &PPCallbacks);