From a1956f64e62dcfdc0e09d286b7f6da1d744254d1 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Wed, 18 Oct 2017 16:26:13 -0400 Subject: [PATCH 01/32] hipDeviceReset(): make sure to reinitialize the printf buffer in hcc RT --- hipamd/src/hip_hcc.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index 2c98ac804f..2d67c31fe7 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -614,6 +614,11 @@ void ihipDevice_t::locked_reset() _state = 0; am_memtracker_reset(_acc); + //FIXME - Calling am_memtracker_reset is really bad since it destroyed all buffers allocated by the HCC runtime as well + //such as the printf buffer. Re-initialze the printf buffer as a workaround for now. +#if (__hcc_workweek__ >= 17423) + Kalmar::getContext()->initPrintfBuffer(); +#endif }; #define ErrorCheck(x) error_check(x, __LINE__, __FILE__) From 72c9767b1dd61d4eef569ba4f866ea23ab82c8b9 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Mon, 16 Oct 2017 20:16:11 +0100 Subject: [PATCH 02/32] Fix broken indentation introduced by previous commit --- hipamd/hipify-clang/CMakeLists.txt | 82 +++++++++++++++--------------- 1 file changed, 41 insertions(+), 41 deletions(-) diff --git a/hipamd/hipify-clang/CMakeLists.txt b/hipamd/hipify-clang/CMakeLists.txt index 1aaa68097f..8bb366a661 100644 --- a/hipamd/hipify-clang/CMakeLists.txt +++ b/hipamd/hipify-clang/CMakeLists.txt @@ -13,6 +13,7 @@ if (NOT ${LLVM_FOUND}) message(STATUS "hipify-clang will not be built. To build it please specify absolute path to LLVM 3.8 or LLVM 3.9 package/dist using -DHIPIFY_CLANG_LLVM_DIR") endif() endif() + if (${LLVM_FOUND}) list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR}) include(AddLLVM) @@ -52,55 +53,54 @@ if (${LLVM_FOUND}) LLVMOption LLVMCore) -if(WIN32) - target_link_libraries(hipify-clang version) -endif() + if(WIN32) + target_link_libraries(hipify-clang version) + endif() set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}") set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS}") -if(MSVC) - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /GR- /EHs- /EHc-") - set(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} /SUBSYSTEM:WINDOWS") -else() - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -pthread -fno-rtti -fvisibility-inlines-hidden") -endif() + if(MSVC) + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /GR- /EHs- /EHc-") + set(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} /SUBSYSTEM:WINDOWS") + else() + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -pthread -fno-rtti -fvisibility-inlines-hidden") + endif() set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DIRS}/clang/${LLVM_VERSION_MAJOR}.${LLVM_VERSION_MINOR}.${LLVM_VERSION_PATCH}\\\"") install(TARGETS hipify-clang DESTINATION bin) -if (HIPIFY_CLANG_TESTS) - # tests - 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.") + if (HIPIFY_CLANG_TESTS) + # tests + 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() + + find_program(LIT_COMMAND lit) + + set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) + + configure_file( + ${CMAKE_SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in + ${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg + @ONLY) + + add_lit_testsuite(test-hipify "Running HIPify regression tests" + ${CMAKE_SOURCE_DIR}/tests/hipify-clang + PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg + DEPENDS hipify-clang lit + ) + + add_custom_target(test-hipify-clang) + add_dependencies(test-hipify-clang test-hipify) + set_target_properties(test-hipify-clang PROPERTIES FOLDER "Tests") endif() - find_program(LIT_COMMAND lit) - - set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) - - configure_file( - ${CMAKE_SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in - ${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg - @ONLY) - - add_lit_testsuite(test-hipify "Running HIPify regression tests" - ${CMAKE_SOURCE_DIR}/tests/hipify-clang - PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg - DEPENDS hipify-clang lit - ) - - add_custom_target(test-hipify-clang) - add_dependencies(test-hipify-clang test-hipify) - set_target_properties(test-hipify-clang PROPERTIES FOLDER "Tests") -endif() - -if (PARENT_SCOPE) - set(BUILD_HIPIFY_CLANG 1 PARENT_SCOPE) -endif() - + if (PARENT_SCOPE) + set(BUILD_HIPIFY_CLANG 1 PARENT_SCOPE) + endif() endif() From d3a56460512bfcbf5289ceb08ede50355aa633df Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Mon, 16 Oct 2017 20:17:29 +0100 Subject: [PATCH 03/32] Use early return to avoid indenting all of CMakeLists.txt --- hipamd/hipify-clang/CMakeLists.txt | 167 ++++++++++++++--------------- 1 file changed, 83 insertions(+), 84 deletions(-) diff --git a/hipamd/hipify-clang/CMakeLists.txt b/hipamd/hipify-clang/CMakeLists.txt index 8bb366a661..7cfac3301a 100644 --- a/hipamd/hipify-clang/CMakeLists.txt +++ b/hipamd/hipify-clang/CMakeLists.txt @@ -11,96 +11,95 @@ if (NOT ${LLVM_FOUND}) find_package(LLVM 3.9 QUIET PATHS ${HIPIFY_CLANG_LLVM_DIR} NO_DEFAULT_PATH) if (NOT ${LLVM_FOUND}) message(STATUS "hipify-clang will not be built. To build it please specify absolute path to LLVM 3.8 or LLVM 3.9 package/dist using -DHIPIFY_CLANG_LLVM_DIR") + return() endif() endif() -if (${LLVM_FOUND}) - list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR}) - include(AddLLVM) +list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR}) +include(AddLLVM) - message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}") +message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}") - include_directories(${LLVM_INCLUDE_DIRS}) - link_directories(${LLVM_LIBRARY_DIRS}) - add_definitions(${LLVM_DEFINITIONS}) - add_llvm_executable(hipify-clang src/Cuda2Hip.cpp) +include_directories(${LLVM_INCLUDE_DIRS}) +link_directories(${LLVM_LIBRARY_DIRS}) +add_definitions(${LLVM_DEFINITIONS}) +add_llvm_executable(hipify-clang src/Cuda2Hip.cpp) - set(CMAKE_CXX_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang++) - set(CMAKE_C_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang) +set(CMAKE_CXX_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang++) +set(CMAKE_C_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang) - # Link against LLVM and CLANG libraries - target_link_libraries(hipify-clang - clangASTMatchers - clangFrontend - clangTooling - clangParse - clangSerialization - clangSema - clangEdit - clangFormat - clangLex - clangAnalysis - clangDriver - clangAST - clangToolingCore - clangRewrite - clangBasic - LLVMProfileData - LLVMSupport - LLVMMCParser - LLVMMC - LLVMBitReader - LLVMOption - LLVMCore) +# Link against LLVM and CLANG libraries +target_link_libraries(hipify-clang + clangASTMatchers + clangFrontend + clangTooling + clangParse + clangSerialization + clangSema + clangEdit + clangFormat + clangLex + clangAnalysis + clangDriver + clangAST + clangToolingCore + clangRewrite + clangBasic + LLVMProfileData + LLVMSupport + LLVMMCParser + LLVMMC + LLVMBitReader + LLVMOption + LLVMCore) - if(WIN32) - target_link_libraries(hipify-clang version) - endif() - - set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}") - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS}") - if(MSVC) - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /GR- /EHs- /EHc-") - set(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} /SUBSYSTEM:WINDOWS") - else() - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -pthread -fno-rtti -fvisibility-inlines-hidden") - endif() - - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DIRS}/clang/${LLVM_VERSION_MAJOR}.${LLVM_VERSION_MINOR}.${LLVM_VERSION_PATCH}\\\"") - - install(TARGETS hipify-clang DESTINATION bin) - - if (HIPIFY_CLANG_TESTS) - # tests - 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() - - find_program(LIT_COMMAND lit) - - set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) - - configure_file( - ${CMAKE_SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in - ${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg - @ONLY) - - add_lit_testsuite(test-hipify "Running HIPify regression tests" - ${CMAKE_SOURCE_DIR}/tests/hipify-clang - PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg - DEPENDS hipify-clang lit - ) - - add_custom_target(test-hipify-clang) - add_dependencies(test-hipify-clang test-hipify) - set_target_properties(test-hipify-clang PROPERTIES FOLDER "Tests") - endif() - - if (PARENT_SCOPE) - set(BUILD_HIPIFY_CLANG 1 PARENT_SCOPE) - endif() +if(WIN32) + target_link_libraries(hipify-clang version) +endif() + +set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS}") +if(MSVC) + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /GR- /EHs- /EHc-") + set(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} /SUBSYSTEM:WINDOWS") +else() + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -pthread -fno-rtti -fvisibility-inlines-hidden") +endif() + +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DIRS}/clang/${LLVM_VERSION_MAJOR}.${LLVM_VERSION_MINOR}.${LLVM_VERSION_PATCH}\\\"") + +install(TARGETS hipify-clang DESTINATION bin) + +if (HIPIFY_CLANG_TESTS) + # tests + 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() + + find_program(LIT_COMMAND lit) + + set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) + + configure_file( + ${CMAKE_SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in + ${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg + @ONLY) + + add_lit_testsuite(test-hipify "Running HIPify regression tests" + ${CMAKE_SOURCE_DIR}/tests/hipify-clang + PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg + DEPENDS hipify-clang lit + ) + + add_custom_target(test-hipify-clang) + add_dependencies(test-hipify-clang test-hipify) + set_target_properties(test-hipify-clang PROPERTIES FOLDER "Tests") +endif() + +if (PARENT_SCOPE) + set(BUILD_HIPIFY_CLANG 1 PARENT_SCOPE) endif() From b926cdf0028ea67c50cb9a9b44200ba5ae92a69a Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Mon, 16 Oct 2017 20:18:06 +0100 Subject: [PATCH 04/32] Declare HIPIFY_CLANG_TESTS as a cmake option --- hipamd/hipify-clang/CMakeLists.txt | 2 ++ 1 file changed, 2 insertions(+) diff --git a/hipamd/hipify-clang/CMakeLists.txt b/hipamd/hipify-clang/CMakeLists.txt index 7cfac3301a..21efb5bc80 100644 --- a/hipamd/hipify-clang/CMakeLists.txt +++ b/hipamd/hipify-clang/CMakeLists.txt @@ -1,6 +1,8 @@ cmake_minimum_required(VERSION 2.8.8) project(hipify-clang) +option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" ON) + if (PARENT_SCOPE) set(BUILD_HIPIFY_CLANG 0 PARENT_SCOPE) endif() From 03dd8c9fcad558f6b20c661de59ae9083d10eec2 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Mon, 16 Oct 2017 20:21:47 +0100 Subject: [PATCH 05/32] Don't be picky about clang versions --- hipamd/hipify-clang/CMakeLists.txt | 10 +--------- 1 file changed, 1 insertion(+), 9 deletions(-) diff --git a/hipamd/hipify-clang/CMakeLists.txt b/hipamd/hipify-clang/CMakeLists.txt index 21efb5bc80..6e3f4e80de 100644 --- a/hipamd/hipify-clang/CMakeLists.txt +++ b/hipamd/hipify-clang/CMakeLists.txt @@ -7,15 +7,7 @@ if (PARENT_SCOPE) set(BUILD_HIPIFY_CLANG 0 PARENT_SCOPE) endif() -# Find LLVM package -find_package(LLVM 3.8 QUIET PATHS ${HIPIFY_CLANG_LLVM_DIR} NO_DEFAULT_PATH) -if (NOT ${LLVM_FOUND}) - find_package(LLVM 3.9 QUIET PATHS ${HIPIFY_CLANG_LLVM_DIR} NO_DEFAULT_PATH) - if (NOT ${LLVM_FOUND}) - message(STATUS "hipify-clang will not be built. To build it please specify absolute path to LLVM 3.8 or LLVM 3.9 package/dist using -DHIPIFY_CLANG_LLVM_DIR") - return() - endif() -endif() +find_package(LLVM PATHS ${HIPIFY_CLANG_LLVM_DIR} REQUIRED) list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR}) include(AddLLVM) From c62766f880eae1d423c08bacc43e7c9d9c18bca1 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Mon, 16 Oct 2017 20:22:12 +0100 Subject: [PATCH 06/32] Don't reinvent find_package PythonInterp is a finder module that ships with cmake. It supports the conventional interaction with find_package that allows you to demand success, and particular vesions, without having your own logic: https://cmake.org/cmake/help/v3.0/command/find_package.html --- hipamd/hipify-clang/CMakeLists.txt | 9 +-------- 1 file changed, 1 insertion(+), 8 deletions(-) diff --git a/hipamd/hipify-clang/CMakeLists.txt b/hipamd/hipify-clang/CMakeLists.txt index 6e3f4e80de..4c6dc8aec8 100644 --- a/hipamd/hipify-clang/CMakeLists.txt +++ b/hipamd/hipify-clang/CMakeLists.txt @@ -66,14 +66,7 @@ install(TARGETS hipify-clang DESTINATION bin) if (HIPIFY_CLANG_TESTS) # tests - 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() - + find_package(PythonInterp 2.7 REQUIRED EXACT) find_program(LIT_COMMAND lit) set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) From 6abd8963aed63fb89217da086cfefb99f559eed7 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Mon, 16 Oct 2017 20:23:32 +0100 Subject: [PATCH 07/32] Skip `lit` tests if `lit` cannot be found --- hipamd/hipify-clang/CMakeLists.txt | 13 ++++++++++--- 1 file changed, 10 insertions(+), 3 deletions(-) diff --git a/hipamd/hipify-clang/CMakeLists.txt b/hipamd/hipify-clang/CMakeLists.txt index 4c6dc8aec8..718e3d15de 100644 --- a/hipamd/hipify-clang/CMakeLists.txt +++ b/hipamd/hipify-clang/CMakeLists.txt @@ -3,6 +3,13 @@ project(hipify-clang) option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" ON) +# Disable the tests if `lit` is not installed. +find_program(LIT_COMMAND lit) +if (NOT LIT_COMMAND) + set(HIPIFY_CLANG_TESTS OFF CACHE INTERNAL "") + message(STATUS "hipify-clang's tests are not being built because `lit` is not installed.") +endif() + if (PARENT_SCOPE) set(BUILD_HIPIFY_CLANG 0 PARENT_SCOPE) endif() @@ -67,20 +74,20 @@ install(TARGETS hipify-clang DESTINATION bin) if (HIPIFY_CLANG_TESTS) # tests find_package(PythonInterp 2.7 REQUIRED EXACT) - find_program(LIT_COMMAND lit) set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) configure_file( ${CMAKE_SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in ${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg - @ONLY) + @ONLY + ) add_lit_testsuite(test-hipify "Running HIPify regression tests" ${CMAKE_SOURCE_DIR}/tests/hipify-clang PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg DEPENDS hipify-clang lit - ) + ) add_custom_target(test-hipify-clang) add_dependencies(test-hipify-clang test-hipify) From ed6681f5635dd61def669c47c22dfa8ec0c07c56 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Mon, 16 Oct 2017 20:24:35 +0100 Subject: [PATCH 08/32] Use the cache for global variables - not PARENT_SCOPE hacks --- hipamd/hipify-clang/CMakeLists.txt | 8 ++------ 1 file changed, 2 insertions(+), 6 deletions(-) diff --git a/hipamd/hipify-clang/CMakeLists.txt b/hipamd/hipify-clang/CMakeLists.txt index 718e3d15de..355b0cea7a 100644 --- a/hipamd/hipify-clang/CMakeLists.txt +++ b/hipamd/hipify-clang/CMakeLists.txt @@ -10,9 +10,7 @@ if (NOT LIT_COMMAND) message(STATUS "hipify-clang's tests are not being built because `lit` is not installed.") endif() -if (PARENT_SCOPE) - set(BUILD_HIPIFY_CLANG 0 PARENT_SCOPE) -endif() +set(BUILD_HIPIFY_CLANG 0 CACHE INTERNAL "") find_package(LLVM PATHS ${HIPIFY_CLANG_LLVM_DIR} REQUIRED) @@ -94,6 +92,4 @@ if (HIPIFY_CLANG_TESTS) set_target_properties(test-hipify-clang PROPERTIES FOLDER "Tests") endif() -if (PARENT_SCOPE) - set(BUILD_HIPIFY_CLANG 1 PARENT_SCOPE) -endif() +set(BUILD_HIPIFY_CLANG 1 CACHE INTERNAL "") From c876f6ffd5379c5c3bc24257db99ac392390224f Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 10:09:48 +0100 Subject: [PATCH 09/32] Fix two faulty LLVM version checks What we actually want to do here is use the StringRef version in versions newer than 3.8, and the void one in 3.8 and older. Checking "major-version >= 3 && minor-version >= 9" does not do what we want. Consider what this will do for version 4.0, for which minor-version is zero... --- hipamd/hipify-clang/src/Cuda2Hip.cpp | 16 ++++++++++------ 1 file changed, 10 insertions(+), 6 deletions(-) diff --git a/hipamd/hipify-clang/src/Cuda2Hip.cpp b/hipamd/hipify-clang/src/Cuda2Hip.cpp index 907e78e54f..5e03e9732d 100644 --- a/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -3109,10 +3109,10 @@ public: // to workaround the 'const' MacroArgs passed into this hook. const Token *start = Args->getUnexpArgument(i); size_t len = Args->getArgLength(start) + 1; -#if (LLVM_VERSION_MAJOR >= 3) && (LLVM_VERSION_MINOR >= 9) - _pp->EnterTokenStream(ArrayRef(start, len), false); -#else +#if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8) _pp->EnterTokenStream(start, len, false, false); +#else + _pp->EnterTokenStream(ArrayRef(start, len), false); #endif do { toks.push_back(Token()); @@ -4207,11 +4207,15 @@ void copyFile(const std::string& src, const std::string& dst) { int main(int argc, const char **argv) { auto start = std::chrono::steady_clock::now(); auto begin = start; -#if (LLVM_VERSION_MAJOR >= 3) && (LLVM_VERSION_MINOR >= 9) - llvm::sys::PrintStackTraceOnErrorSignal(StringRef()); -#else + + // The signature of PrintStackTraceOnErrorSignal changed in llvm 3.9. We don't support + // anything older than 3.8, so let's specifically detect the one old version we support. +#if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8) llvm::sys::PrintStackTraceOnErrorSignal(); +#else + llvm::sys::PrintStackTraceOnErrorSignal(StringRef()); #endif + CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::OneOrMore); std::vector fileSources = OptionsParser.getSourcePathList(); std::string dst = OutputFilename; From 73984ed8094db75070a37c7120cf7c4cc8b57081 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 10:16:39 +0100 Subject: [PATCH 10/32] Refer to clang::StringLiteral explicitly Newer versions of llvm/clang mean there is both an llvm::StringLiteral and a clang::StringLiteral. Since we're dumping both namespaces wholesale into the global namespace with `using` declarations, this creates a name collision, which must be resolved. This change is backwards-compatible, and fixes a problem you encounter when using newer versions of the llvm/clang API. --- hipamd/hipify-clang/src/Cuda2Hip.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/hipify-clang/src/Cuda2Hip.cpp b/hipamd/hipify-clang/src/Cuda2Hip.cpp index 5e03e9732d..cf11dc9546 100644 --- a/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -3799,7 +3799,7 @@ private: } bool stringLiteral(const MatchFinder::MatchResult &Result) { - if (const StringLiteral *sLiteral = Result.Nodes.getNodeAs("stringLiteral")) { + if (const clang::StringLiteral *sLiteral = Result.Nodes.getNodeAs("stringLiteral")) { if (sLiteral->getCharByteWidth() == 1) { StringRef s = sLiteral->getString(); SourceManager *SM = Result.SourceManager; From 22a5e2330d94257cd45c8b611b641aab71427ea5 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 10:29:47 +0100 Subject: [PATCH 11/32] Remove unnecessary call to Retain() on a smart pointer The Preprocessor smart pointer is held by the CompilerInstance, and therefore its reference count cannot reach zero until the CompilerInstance itself is destroyed. If the CompilerInstance is destroyed, you have more to worry about than just the preprocessor being deallocated! Newer versions of the LLVM/Clang API migrated to using std::shared_ptr, so there is no `Retain()` function (by that name, anyway). Eliminating this redundant use is a neat and backward-compatible way to become compatible with newer versions of the LLVM/Clang API. --- hipamd/hipify-clang/src/Cuda2Hip.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/hipamd/hipify-clang/src/Cuda2Hip.cpp b/hipamd/hipify-clang/src/Cuda2Hip.cpp index cf11dc9546..fc224adaf0 100644 --- a/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -3029,7 +3029,6 @@ public: SourceManager &SM = CI.getSourceManager(); setSourceManager(&SM); PP.addPPCallbacks(std::unique_ptr(this)); - PP.Retain(); setPreprocessor(&PP); return true; } From 9ccb966e2bc02c1745fd8d1b521efdc5490ebfe1 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 10:38:24 +0100 Subject: [PATCH 12/32] Cope with clang 4.0's rename of getNumArgs() Sorry, this one I couldn't do in a perfectly elegant way ;) --- hipamd/hipify-clang/src/Cuda2Hip.cpp | 9 ++++++++- 1 file changed, 8 insertions(+), 1 deletion(-) diff --git a/hipamd/hipify-clang/src/Cuda2Hip.cpp b/hipamd/hipify-clang/src/Cuda2Hip.cpp index fc224adaf0..ef53afaaf3 100644 --- a/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -3102,7 +3102,14 @@ public: const MacroDefinition &MD, SourceRange Range, const MacroArgs *Args) override { if (_sm->isWrittenInMainFile(MacroNameTok.getLocation())) { - for (unsigned int i = 0; Args && i < MD.getMacroInfo()->getNumArgs(); i++) { + // The getNumArgs function was rather unhelpfully renamed in clang 4.0. Its semantics + // remain unchanged. +#if LLVM_VERSION_MAJOR > 3 + #define GET_NUM_ARGS() getNumParams() +#else + #define GET_NUM_ARGS() getNumArgs() +#endif + for (unsigned int i = 0; Args && i < MD.getMacroInfo()->GET_NUM_ARGS(); i++) { std::vector toks; // Code below is a kind of stolen from 'MacroArgs::getPreExpArgument' // to workaround the 'const' MacroArgs passed into this hook. From 885573e1e2fba3db76d8a7d101eb6a16806dcc4c Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 10:56:31 +0100 Subject: [PATCH 13/32] Cope with Replacements now having llvm::Error returns --- hipamd/hipify-clang/src/Cuda2Hip.cpp | 6 ++++++ 1 file changed, 6 insertions(+) diff --git a/hipamd/hipify-clang/src/Cuda2Hip.cpp b/hipamd/hipify-clang/src/Cuda2Hip.cpp index ef53afaaf3..816ca4bc3d 100644 --- a/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -2928,7 +2928,13 @@ protected: std::string mainFileName; virtual void insertReplacement(const Replacement &rep, const FullSourceLoc &fullSL) { +#if LLVM_VERSION_MAJOR > 3 + // New clang added error checking to Replacements, and *insists* that you explicitly check it. + llvm::Error e = Replace->add(rep); +#else + // In older versions, it's literally an std::set Replace->insert(rep); +#endif if (PrintStats) { LOCs.insert(fullSL.getExpansionLineNumber()); } From abd34b2125403f62d5778a9446445eb10304cd67 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 10:56:46 +0100 Subject: [PATCH 14/32] Omit now-removed Filename string arg from handleBeginSource --- hipamd/hipify-clang/src/Cuda2Hip.cpp | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/hipamd/hipify-clang/src/Cuda2Hip.cpp b/hipamd/hipify-clang/src/Cuda2Hip.cpp index 816ca4bc3d..0f095c7245 100644 --- a/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -3030,7 +3030,11 @@ public: HipifyPPCallbacks(Replacements *R, const std::string &mainFileName) : Cuda2Hip(R, mainFileName), SeenEnd(false), _sm(nullptr), _pp(nullptr) {} - virtual bool handleBeginSource(CompilerInstance &CI, StringRef Filename) override { + virtual bool handleBeginSource(CompilerInstance &CI +#if LLVM_VERSION_MAJOR < 4 + , StringRef Filename +#endif + ) override { Preprocessor &PP = CI.getPreprocessor(); SourceManager &SM = CI.getSourceManager(); setSourceManager(&SM); From 9d83348ac72addf50a3487fe2c33a1a5d9d37322 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 11:09:17 +0100 Subject: [PATCH 15/32] Use inline initialisers to set default field values A trivial cleanup that helps in a moment.. --- hipamd/hipify-clang/src/Cuda2Hip.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/hipamd/hipify-clang/src/Cuda2Hip.cpp b/hipamd/hipify-clang/src/Cuda2Hip.cpp index 0f095c7245..b71a55d81d 100644 --- a/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -3028,7 +3028,7 @@ class Cuda2HipCallback; class HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks, public Cuda2Hip { public: HipifyPPCallbacks(Replacements *R, const std::string &mainFileName) - : Cuda2Hip(R, mainFileName), SeenEnd(false), _sm(nullptr), _pp(nullptr) {} + : Cuda2Hip(R, mainFileName) {} virtual bool handleBeginSource(CompilerInstance &CI #if LLVM_VERSION_MAJOR < 4 @@ -3201,15 +3201,15 @@ public: void EndOfMainFile() override {} - bool SeenEnd; + bool SeenEnd = false; void setSourceManager(SourceManager *sm) { _sm = sm; } void setPreprocessor(Preprocessor *pp) { _pp = pp; } void setMatch(Cuda2HipCallback *match) { Match = match; } private: - SourceManager *_sm; - Preprocessor *_pp; - Cuda2HipCallback *Match; + SourceManager *_sm = nullptr; + Preprocessor *_pp = nullptr; + Cuda2HipCallback *Match = nullptr; }; class Cuda2HipCallback : public MatchFinder::MatchCallback, public Cuda2Hip { From a06fe54ee8027a6530424c6841b70618846b0bd3 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 11:43:35 +0100 Subject: [PATCH 16/32] Be agnostic to the new getReplacements() API See comment --- hipamd/hipify-clang/src/Cuda2Hip.cpp | 26 +++++++++++++++++++++----- 1 file changed, 21 insertions(+), 5 deletions(-) diff --git a/hipamd/hipify-clang/src/Cuda2Hip.cpp b/hipamd/hipify-clang/src/Cuda2Hip.cpp index b71a55d81d..8b16e75db0 100644 --- a/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -4291,8 +4291,19 @@ int main(int argc, const char **argv) { // So what we do is operate on a copy, which we then move to the output. RefactoringTool Tool(OptionsParser.getCompilations(), tmpFile); ast_matchers::MatchFinder Finder; - HipifyPPCallbacks PPCallbacks(&Tool.getReplacements(), tmpFile); - Cuda2HipCallback Callback(&Tool.getReplacements(), &Finder, &PPCallbacks, tmpFile); + + // The Replacements to apply to the file `src`. + Replacements* replacementsToUse; +#if LLVM_VERSION_MAJOR > 3 + // getReplacements() now returns a map from filename to Replacements - so create an entry + // for this source file and return a pointer to it. + replacementsToUse = &(Tool.getReplacements()[tmpFile]); +#else + replacementsToUse = &Tool.getReplacements(); +#endif + + HipifyPPCallbacks PPCallbacks(replacementsToUse, tmpFile); + Cuda2HipCallback Callback(replacementsToUse, &Finder, &PPCallbacks, tmpFile); addAllMatchers(Finder, &Callback); @@ -4320,9 +4331,14 @@ int main(int argc, const char **argv) { SourceManager SM(Diagnostics, Tool.getFiles()); if (PrintStats) { DEBUG(dbgs() << "Replacements collected by the tool:\n"); - for (const auto &r : Tool.getReplacements()) { - DEBUG(dbgs() << r.toString() << "\n"); - repBytes += r.getLength(); +#if LLVM_VERSION_MAJOR > 3 + Replacements& replacements = Tool.getReplacements().begin()->second; +#else + Replacements& replacements = Tool.getReplacements(); +#endif + for (const auto &replacement : replacements) { + DEBUG(dbgs() << replacement.toString() << "\n"); + repBytes += replacement.getLength(); } std::ifstream src_file(dst, std::ios::binary | std::ios::ate); src_file.clear(); From 893ee6d6ca12f353f6b16846c39411b587e552dc Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Sun, 15 Oct 2017 11:51:35 +0100 Subject: [PATCH 17/32] Avoid a double-free of HipifyPPCallbacks instance This bug was present all along, but something changed in the order of de-initialisation performed by llvm that makes it actually crash now. The constructor of HipifyPPCallbacks gives: ``` std::unique_ptr(this) ``` to the LLVM Preprocessor instance. The Preprocessor instance subsequently frees the HipifyPPCallbacks, which is then freed again when we leave the stack frame at line 4340. So: let's leak the HipifyPPCallbacks onto the heap, and leave the LLVM Preprocessor object responsible for tidying it up. --- hipamd/hipify-clang/src/Cuda2Hip.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/hipamd/hipify-clang/src/Cuda2Hip.cpp b/hipamd/hipify-clang/src/Cuda2Hip.cpp index 8b16e75db0..0f73859893 100644 --- a/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -4302,12 +4302,12 @@ int main(int argc, const char **argv) { replacementsToUse = &Tool.getReplacements(); #endif - HipifyPPCallbacks PPCallbacks(replacementsToUse, tmpFile); - Cuda2HipCallback Callback(replacementsToUse, &Finder, &PPCallbacks, tmpFile); + HipifyPPCallbacks* PPCallbacks = new HipifyPPCallbacks(replacementsToUse, tmpFile); + Cuda2HipCallback Callback(replacementsToUse, &Finder, PPCallbacks, tmpFile); addAllMatchers(Finder, &Callback); - auto action = newFrontendActionFactory(&Finder, &PPCallbacks); + auto action = newFrontendActionFactory(&Finder, PPCallbacks); Tool.appendArgumentsAdjuster(getInsertArgumentAdjuster("--cuda-host-only", ArgumentInsertPosition::BEGIN)); @@ -4365,13 +4365,13 @@ int main(int argc, const char **argv) { } std::remove(csv.c_str()); } - if (0 == printStats(csv, src, PPCallbacks, Callback, repBytes, bytes, lines, start)) { + if (0 == printStats(csv, src, *PPCallbacks, Callback, repBytes, bytes, lines, start)) { filesTranslated--; } start = std::chrono::steady_clock::now(); repBytesTotal += repBytes; bytesTotal += bytes; - changedLinesTotal += PPCallbacks.LOCs.size() + Callback.LOCs.size(); + changedLinesTotal += PPCallbacks->LOCs.size() + Callback.LOCs.size(); linesTotal += lines; } dst.clear(); From e9d259699c96d723e8cd13a07cedb69a2972692f Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 08:00:58 +0100 Subject: [PATCH 18/32] Remove dependency on nonexistent `lit` target This just... isn't a target. Maybe you were trying to do a file dependency? But that's handled implicitly by the find_program call anyway, so all this is doing is generating a cmake warning --- hipamd/hipify-clang/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/hipify-clang/CMakeLists.txt b/hipamd/hipify-clang/CMakeLists.txt index 355b0cea7a..d309aa261c 100644 --- a/hipamd/hipify-clang/CMakeLists.txt +++ b/hipamd/hipify-clang/CMakeLists.txt @@ -84,7 +84,7 @@ if (HIPIFY_CLANG_TESTS) add_lit_testsuite(test-hipify "Running HIPify regression tests" ${CMAKE_SOURCE_DIR}/tests/hipify-clang PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg - DEPENDS hipify-clang lit + DEPENDS hipify-clang ) add_custom_target(test-hipify-clang) From 9747578d093164d7e5325aa1ef79958879e17b35 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 08:16:26 +0100 Subject: [PATCH 19/32] Propagate the CUDA toolkit directory into the lit tests Allows the tests to actually run... :D --- hipamd/hipify-clang/CMakeLists.txt | 5 ++++- hipamd/tests/hipify-clang/axpy.cu | 2 +- hipamd/tests/hipify-clang/lit.cfg | 2 ++ hipamd/tests/hipify-clang/lit.site.cfg.in | 1 + 4 files changed, 8 insertions(+), 2 deletions(-) diff --git a/hipamd/hipify-clang/CMakeLists.txt b/hipamd/hipify-clang/CMakeLists.txt index d309aa261c..18939c0210 100644 --- a/hipamd/hipify-clang/CMakeLists.txt +++ b/hipamd/hipify-clang/CMakeLists.txt @@ -70,9 +70,12 @@ set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DI install(TARGETS hipify-clang DESTINATION bin) if (HIPIFY_CLANG_TESTS) - # tests find_package(PythonInterp 2.7 REQUIRED EXACT) + # Populates CUDA_TOOLKIT_ROOT_DIR, which is then applied to the lit config to give the + # value of --cuda-path for the test runs. + find_package(CUDA REQUIRED) + set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) configure_file( diff --git a/hipamd/tests/hipify-clang/axpy.cu b/hipamd/tests/hipify-clang/axpy.cu index 8472e60209..56a854fd0a 100644 --- a/hipamd/tests/hipify-clang/axpy.cu +++ b/hipamd/tests/hipify-clang/axpy.cu @@ -1,4 +1,4 @@ -// RUN: hipify "%s" -o=%t -- +// RUN: hipify "%s" -o=%t -- %cuda_args #include diff --git a/hipamd/tests/hipify-clang/lit.cfg b/hipamd/tests/hipify-clang/lit.cfg index c57b8ec524..57d9d876b5 100644 --- a/hipamd/tests/hipify-clang/lit.cfg +++ b/hipamd/tests/hipify-clang/lit.cfg @@ -46,3 +46,5 @@ if obj_root is not None: config.substitutions.append(("hipify", obj_root+"/hipify-clang")) +# Clang args for CUDA... +config.substitutions.append(("%cuda_args", "-x cuda --cuda-path=%s --cuda-gpu-arch=sm_30" % config.cuda_root)) diff --git a/hipamd/tests/hipify-clang/lit.site.cfg.in b/hipamd/tests/hipify-clang/lit.site.cfg.in index 4511316ac7..4e05710065 100644 --- a/hipamd/tests/hipify-clang/lit.site.cfg.in +++ b/hipamd/tests/hipify-clang/lit.site.cfg.in @@ -2,6 +2,7 @@ import sys config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@" config.obj_root = "@BINARY_DIR@" +config.cuda_root = "@CUDA_TOOLKIT_ROOT_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. From 3868036ea7c80f46e9f305775d1145b883b5d0bb Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 08:27:22 +0100 Subject: [PATCH 20/32] Look for FileCheck for running lit tests, too Use of grep in `lit` RUN lines is deprecated: https://llvm.org/docs/TestingGuide.html#writing-new-regression-tests Using grep leads to really unhelpful failure output (it literally just says "the test failed"). FileCheck is much more helpful, and distributed with LLVM on most distros anyway, so this extra dependency shouldn't prove problematic. --- hipamd/hipify-clang/CMakeLists.txt | 7 ++++--- 1 file changed, 4 insertions(+), 3 deletions(-) diff --git a/hipamd/hipify-clang/CMakeLists.txt b/hipamd/hipify-clang/CMakeLists.txt index 18939c0210..e7c46557f8 100644 --- a/hipamd/hipify-clang/CMakeLists.txt +++ b/hipamd/hipify-clang/CMakeLists.txt @@ -3,11 +3,12 @@ project(hipify-clang) option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" ON) -# Disable the tests if `lit` is not installed. +# Disable the tests if `lit` or `FileCheck` is not installed. find_program(LIT_COMMAND lit) -if (NOT LIT_COMMAND) +find_program(FILECHECK_COMMAND FileCheck) +if (NOT LIT_COMMAND OR NOT FILECHECK_COMMAND) set(HIPIFY_CLANG_TESTS OFF CACHE INTERNAL "") - message(STATUS "hipify-clang's tests are not being built because `lit` is not installed.") + message(STATUS "hipify-clang's tests are not being built because `lit` or `FileCheck` could not be found.") endif() set(BUILD_HIPIFY_CLANG 0 CACHE INTERNAL "") From 74fd64d5c132081b56301e07b1220a608d8d8ac0 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 08:45:16 +0100 Subject: [PATCH 21/32] Migrate lit test to using FileCheck, so failures are readable It seems the test is already broken, but look how awesome the error message is now: /home/chris/HIP/tests/hipify-clang/axpy.cu:31:12: error: expected string not found in input // CHECK: hipLaunchKernel(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); ^ :31:2: note: scanning from here // ^ :33:2: note: possible intended match here hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); ^ --- hipamd/tests/hipify-clang/axpy.cu | 22 +++++++++++++--------- 1 file changed, 13 insertions(+), 9 deletions(-) diff --git a/hipamd/tests/hipify-clang/axpy.cu b/hipamd/tests/hipify-clang/axpy.cu index 56a854fd0a..e0afa02a26 100644 --- a/hipamd/tests/hipify-clang/axpy.cu +++ b/hipamd/tests/hipify-clang/axpy.cu @@ -1,9 +1,9 @@ -// RUN: hipify "%s" -o=%t -- %cuda_args +// RUN: hipify "%s" -o=%t -- %cuda_args && cat %t | sed -Ee 's|//.+|// |g' | FileCheck %s #include __global__ void axpy(float a, float* x, float* y) { - // RUN: sh -c "test `grep -c -F 'y[hipThreadIdx_x] = a * x[hipThreadIdx_x];' %t` -eq 2" + // CHECK: y[hipThreadIdx_x] = a * x[hipThreadIdx_x]; y[threadIdx.x] = a * x[threadIdx.x]; } @@ -17,21 +17,25 @@ int main(int argc, char* argv[]) { // Copy input data to device. float* device_x; float* device_y; - // RUN: sh -c "test `grep -c -F 'hipMalloc(&device_x, kDataLen * sizeof(float));' %t` -eq 2" + + // CHECK: hipMalloc(&device_x, kDataLen * sizeof(float)); cudaMalloc(&device_x, kDataLen * sizeof(float)); - // RUN: sh -c "test `grep -c -F 'hipMalloc(&device_y, kDataLen * sizeof(float));' %t` -eq 2" + + // CHECK: hipMalloc(&device_y, kDataLen * sizeof(float)); cudaMalloc(&device_y, kDataLen * sizeof(float)); - // RUN: sh -c "test `grep -c -F 'hipMemcpy(device_x, host_x, kDataLen * sizeof(float), hipMemcpyHostToDevice);' %t` -eq 2" + + // CHECK: hipMemcpy(device_x, host_x, kDataLen * sizeof(float), hipMemcpyHostToDevice); cudaMemcpy(device_x, host_x, kDataLen * sizeof(float), cudaMemcpyHostToDevice); // Launch the kernel. - // RUN: sh -c "test `grep -c -F 'hipLaunchKernel(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);' %t` -eq 2" + // CHECK: hipLaunchKernel(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); axpy<<<1, kDataLen>>>(a, device_x, device_y); // Copy output data to host. - // RUN: sh -c "test `grep -c -F 'hipDeviceSynchronize();' %t` -eq 2" + // CHECK: hipDeviceSynchronize(); cudaDeviceSynchronize(); - // RUN: sh -c "test `grep -c -F 'hipMemcpy(host_y, device_y, kDataLen * sizeof(float), hipMemcpyDeviceToHost);' %t` -eq 2" + + // CHECK: hipMemcpy(host_y, device_y, kDataLen * sizeof(float), hipMemcpyDeviceToHost); cudaMemcpy(host_y, device_y, kDataLen * sizeof(float), cudaMemcpyDeviceToHost); // Print the results. @@ -39,7 +43,7 @@ int main(int argc, char* argv[]) { std::cout << "y[" << i << "] = " << host_y[i] << "\n"; } - // RUN: sh -c "test `grep -c -F 'hipDeviceReset();' %t` -eq 2" + // CHECK: hipDeviceReset(); cudaDeviceReset(); return 0; } From 5912f465bdc9bed19e3e6dbf7872e73de7c60383 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 08:46:36 +0100 Subject: [PATCH 22/32] Adapt `lit` test for the hipLaunchKernelGGL changes from before... --- hipamd/tests/hipify-clang/axpy.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/tests/hipify-clang/axpy.cu b/hipamd/tests/hipify-clang/axpy.cu index e0afa02a26..689c6e96ff 100644 --- a/hipamd/tests/hipify-clang/axpy.cu +++ b/hipamd/tests/hipify-clang/axpy.cu @@ -28,7 +28,7 @@ int main(int argc, char* argv[]) { cudaMemcpy(device_x, host_x, kDataLen * sizeof(float), cudaMemcpyHostToDevice); // Launch the kernel. - // CHECK: hipLaunchKernel(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); + // CHECK: hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); axpy<<<1, kDataLen>>>(a, device_x, device_y); // Copy output data to host. From c99dcbba8d5e31c64098bb707f892d193e4da322 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 13:50:00 +0100 Subject: [PATCH 23/32] Introduce a test runner script to simplify invocation ... And to use a standard, highly amusing trick for making coloured output work. --- hipamd/hipify-clang/CMakeLists.txt | 6 ++++-- hipamd/tests/hipify-clang/axpy.cu | 2 +- hipamd/tests/hipify-clang/lit.cfg | 1 + hipamd/tests/hipify-clang/run_test.sh | 28 +++++++++++++++++++++++++++ 4 files changed, 34 insertions(+), 3 deletions(-) create mode 100755 hipamd/tests/hipify-clang/run_test.sh diff --git a/hipamd/hipify-clang/CMakeLists.txt b/hipamd/hipify-clang/CMakeLists.txt index e7c46557f8..2ec96e51a8 100644 --- a/hipamd/hipify-clang/CMakeLists.txt +++ b/hipamd/hipify-clang/CMakeLists.txt @@ -6,9 +6,10 @@ option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed # Disable the tests if `lit` or `FileCheck` is not installed. find_program(LIT_COMMAND lit) find_program(FILECHECK_COMMAND FileCheck) -if (NOT LIT_COMMAND OR NOT FILECHECK_COMMAND) +find_program(SOCAT_COMMAND socat) +if (NOT LIT_COMMAND OR NOT FILECHECK_COMMAND OR NOT SOCAT_COMMAND) set(HIPIFY_CLANG_TESTS OFF CACHE INTERNAL "") - message(STATUS "hipify-clang's tests are not being built because `lit` or `FileCheck` could not be found.") + message(STATUS "hipify-clang's tests are not being built because `lit`,`FileCheck` or `socat` could not be found.") endif() set(BUILD_HIPIFY_CLANG 0 CACHE INTERNAL "") @@ -88,6 +89,7 @@ if (HIPIFY_CLANG_TESTS) add_lit_testsuite(test-hipify "Running HIPify regression tests" ${CMAKE_SOURCE_DIR}/tests/hipify-clang PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg + ARGS -v DEPENDS hipify-clang ) diff --git a/hipamd/tests/hipify-clang/axpy.cu b/hipamd/tests/hipify-clang/axpy.cu index 689c6e96ff..8c6b0e0d8d 100644 --- a/hipamd/tests/hipify-clang/axpy.cu +++ b/hipamd/tests/hipify-clang/axpy.cu @@ -1,4 +1,4 @@ -// RUN: hipify "%s" -o=%t -- %cuda_args && cat %t | sed -Ee 's|//.+|// |g' | FileCheck %s +// RUN: %run_test hipify "%s" "%t" %cuda_args #include diff --git a/hipamd/tests/hipify-clang/lit.cfg b/hipamd/tests/hipify-clang/lit.cfg index 57d9d876b5..bb6ac2b407 100644 --- a/hipamd/tests/hipify-clang/lit.cfg +++ b/hipamd/tests/hipify-clang/lit.cfg @@ -48,3 +48,4 @@ config.substitutions.append(("hipify", obj_root+"/hipify-clang")) # Clang args for CUDA... config.substitutions.append(("%cuda_args", "-x cuda --cuda-path=%s --cuda-gpu-arch=sm_30" % config.cuda_root)) +config.substitutions.append(("%run_test", config.test_source_root + "/run_test.sh")) diff --git a/hipamd/tests/hipify-clang/run_test.sh b/hipamd/tests/hipify-clang/run_test.sh new file mode 100755 index 0000000000..46b2fc066b --- /dev/null +++ b/hipamd/tests/hipify-clang/run_test.sh @@ -0,0 +1,28 @@ +#!/usr/bin/env bash + +set -o errexit + +# Run a single LIT test file in a magical way that preserves colour output, to work around +# a known flaw in lit. + +# Capture lit substitutions +HIPIFY=$1 +IN_FILE=$2 +TMP_FILE=$3 +shift 3 + +# Remaining args are the ones to forward to clang proper. + +# Time for the classic insane little trick for making colour output work. +# A self-deleting shell-script that does the thing we want to do... +TMP_SCRIPT=$(mktemp) +cat << EOF > $TMP_SCRIPT +set -o errexit +set -o xtrace +rm $TMP_SCRIPT +$HIPIFY -o=$TMP_FILE $IN_FILE -- $@ && cat $TMP_FILE | sed -Ee 's|//.+|// |g' | FileCheck $IN_FILE +EOF +chmod a+x $TMP_SCRIPT + +# Run the script via socat, spawning a virtual terminal and propagating exit code, and hence failure. +socat -du EXEC:$TMP_SCRIPT,pty,stderr STDOUT From ead79e5bf4e54f0f7b1fda3482c2d1b0a36d8022 Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 13:58:34 +0100 Subject: [PATCH 24/32] Add square.cu to lit testsuite --- hipamd/tests/hipify-clang/square.cu | 114 ++++++++++++++++++++++++++++ 1 file changed, 114 insertions(+) create mode 100644 hipamd/tests/hipify-clang/square.cu diff --git a/hipamd/tests/hipify-clang/square.cu b/hipamd/tests/hipify-clang/square.cu new file mode 100644 index 0000000000..e0c72094a8 --- /dev/null +++ b/hipamd/tests/hipify-clang/square.cu @@ -0,0 +1,114 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +/* +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. +*/ +#include +#include + +#define CHECK(cmd) \ +{\ + cudaError_t error = cmd;\ + if (error != cudaSuccess) { \ + fprintf(stderr, "error: '%s'(%d) at %s:%d\n", cudaGetErrorString(error), error,__FILE__, __LINE__); \ + exit(EXIT_FAILURE);\ + }\ +} + + +/* + * Square each element in the array A and write to array C. + */ +template +__global__ void +vector_square(T *C_d, const T *A_d, size_t N) +{ + // CHECK: size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); + // CHECK: size_t stride = hipBlockDim_x * hipGridDim_x; + size_t offset = (blockIdx.x * blockDim.x + threadIdx.x); + size_t stride = blockDim.x * gridDim.x; + + for (size_t i=offset; i>> (C_d, A_d, N); + + printf ("info: copy Device2Host\n"); + // CHECK: CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + CHECK ( cudaMemcpy(C_h, C_d, Nbytes, cudaMemcpyDeviceToHost)); + + printf ("info: check result\n"); + for (size_t i=0; i Date: Wed, 18 Oct 2017 14:05:46 +0100 Subject: [PATCH 25/32] Add cudaRegister.cu lit test --- hipamd/tests/hipify-clang/cudaRegister.cu | 111 ++++++++++++++++++++++ 1 file changed, 111 insertions(+) create mode 100644 hipamd/tests/hipify-clang/cudaRegister.cu diff --git a/hipamd/tests/hipify-clang/cudaRegister.cu b/hipamd/tests/hipify-clang/cudaRegister.cu new file mode 100644 index 0000000000..80d17f65b9 --- /dev/null +++ b/hipamd/tests/hipify-clang/cudaRegister.cu @@ -0,0 +1,111 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +/* +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. +*/ + +#include +#include +#include +#include +#include +#include + +#define LEN 1024 +#define SIZE LEN * sizeof(float) +#define ITER 1024*1024 + +// CHECK: if(status != hipSuccess) { +#define check(msg, status){ \ +if(status != cudaSuccess) { \ + printf("%s failed. \n", #msg); \ +} \ +} + +__global__ void Inc1(float *Ad, float *Bd){ + // CHECK: int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + int tx = threadIdx.x + blockIdx.x * blockDim.x; + if(tx < 1 ){ + for(int i=0;i>>(Ad, Bd); + sleep(3); + A[0] = -(ITER*1.0f); + std::cout<<"Same cache line before completion: \t"<< A[0]<>>(Ad, Bd); + sleep(3); + A[0] = -(ITER*1.0f); + std::cout<<"Diff cache line before completion: \t"< Date: Wed, 18 Oct 2017 20:41:23 +0100 Subject: [PATCH 26/32] Add the CUDA samples include dir to the path for tests Means we get to easily steal CUDA examples for tests --- hipamd/tests/hipify-clang/lit.cfg | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/tests/hipify-clang/lit.cfg b/hipamd/tests/hipify-clang/lit.cfg index bb6ac2b407..76b2ca08dc 100644 --- a/hipamd/tests/hipify-clang/lit.cfg +++ b/hipamd/tests/hipify-clang/lit.cfg @@ -47,5 +47,5 @@ if obj_root is not None: config.substitutions.append(("hipify", obj_root+"/hipify-clang")) # Clang args for CUDA... -config.substitutions.append(("%cuda_args", "-x cuda --cuda-path=%s --cuda-gpu-arch=sm_30" % config.cuda_root)) +config.substitutions.append(("%cuda_args", "-x cuda --cuda-path=%s --cuda-gpu-arch=sm_30 -isystem%s/samples/common/inc" % (config.cuda_root, config.cuda_root))) config.substitutions.append(("%run_test", config.test_source_root + "/run_test.sh")) From 2c65d0da37540e27d67cdd2b357f5d5e5a0aa22e Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Wed, 18 Oct 2017 21:12:35 +0100 Subject: [PATCH 27/32] Add concurentKernels.cu to the testsuite --- hipamd/tests/hipify-clang/concurentKernels.cu | 240 ++++++++++++++++++ 1 file changed, 240 insertions(+) create mode 100644 hipamd/tests/hipify-clang/concurentKernels.cu diff --git a/hipamd/tests/hipify-clang/concurentKernels.cu b/hipamd/tests/hipify-clang/concurentKernels.cu new file mode 100644 index 0000000000..e369baaf3e --- /dev/null +++ b/hipamd/tests/hipify-clang/concurentKernels.cu @@ -0,0 +1,240 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args +/* + * Copyright 1993-2015 NVIDIA Corporation. All rights reserved. + * + * Please refer to the NVIDIA end user license agreement (EULA) associated + * with this source code for terms and conditions that govern your use of + * this software. Any use, reproduction, disclosure, or distribution of + * this software and related documentation outside the terms of the EULA + * is strictly prohibited. + * + */ + +// +// This sample demonstrates the use of streams for concurrent execution. It also illustrates how to +// introduce dependencies between CUDA streams with the new cudaStreamWaitEvent function introduced +// in CUDA 3.2. +// +// Devices of compute capability 1.x will run the kernels one after another +// Devices of compute capability 2.0 or higher can overlap the kernels +// +#include +#include +#include + +// This is a kernel that does no real work but runs at least for a specified number of clocks +__global__ void clock_block(clock_t *d_o, clock_t clock_count) +{ + unsigned int start_clock = (unsigned int) clock(); + + clock_t clock_offset = 0; + + while (clock_offset < clock_count) + { + unsigned int end_clock = (unsigned int) clock(); + + // The code below should work like + // this (thanks to modular arithmetics): + // + // clock_offset = (clock_t) (end_clock > start_clock ? + // end_clock - start_clock : + // end_clock + (0xffffffffu - start_clock)); + // + // Indeed, let m = 2^32 then + // end - start = end + m - start (mod m). + + clock_offset = (clock_t)(end_clock - start_clock); + } + + d_o[0] = clock_offset; +} + + +// Single warp reduction kernel +__global__ void sum(clock_t *d_clocks, int N) +{ + __shared__ clock_t s_clocks[32]; + + clock_t my_sum = 0; + + for (int i = threadIdx.x; i < N; i+= blockDim.x) + { + my_sum += d_clocks[i]; + } + + s_clocks[threadIdx.x] = my_sum; + syncthreads(); + + for (int i=16; i>0; i/=2) + { + if (threadIdx.x < i) + { + s_clocks[threadIdx.x] += s_clocks[threadIdx.x + i]; + } + + syncthreads(); + } + + d_clocks[0] = s_clocks[0]; +} + +int main(int argc, char **argv) +{ + int nkernels = 8; // number of concurrent kernels + int nstreams = nkernels + 1; // use one more stream than concurrent kernel + int nbytes = nkernels * sizeof(clock_t); // number of data bytes + float kernel_time = 10; // time the kernel should run in ms + float elapsed_time; // timing variables + int cuda_device = 0; + + printf("[%s] - Starting...\n", argv[0]); + + // get number of kernels if overridden on the command line + if (checkCmdLineFlag(argc, (const char **)argv, "nkernels")) + { + nkernels = getCmdLineArgumentInt(argc, (const char **)argv, "nkernels"); + nstreams = nkernels + 1; + } + + // use command-line specified CUDA device, otherwise use device with highest Gflops/s + cuda_device = findCudaDevice(argc, (const char **)argv); + + cudaDeviceProp deviceProp; + // CHECK: checkCudaErrors(hipGetDevice(&cuda_device)); + checkCudaErrors(cudaGetDevice(&cuda_device)); + + // CHECK: checkCudaErrors(hipGetDeviceProperties(&deviceProp, cuda_device)); + checkCudaErrors(cudaGetDeviceProperties(&deviceProp, cuda_device)); + + if ((deviceProp.concurrentKernels == 0)) + { + printf("> GPU does not support concurrent kernel execution\n"); + printf(" CUDA kernel runs will be serialized\n"); + } + + printf("> Detected Compute SM %d.%d hardware with %d multi-processors\n", + deviceProp.major, deviceProp.minor, deviceProp.multiProcessorCount); + + // allocate host memory + clock_t *a = 0; // pointer to the array data in host memory + // CHECK: checkCudaErrors(hipHostMalloc((void **)&a, nbytes)); + checkCudaErrors(cudaMallocHost((void **)&a, nbytes)); + + // allocate device memory + clock_t *d_a = 0; // pointers to data and init value in the device memory + // CHECK: checkCudaErrors(hipMalloc((void **)&d_a, nbytes)); + checkCudaErrors(cudaMalloc((void **)&d_a, nbytes)); + + // CHECK: hipStream_t *streams = (hipStream_t *) malloc(nstreams * sizeof(hipStream_t)); + // allocate and initialize an array of stream handles + cudaStream_t *streams = (cudaStream_t *) malloc(nstreams * sizeof(cudaStream_t)); + + for (int i = 0; i < nstreams; i++) + { + // CHECK: checkCudaErrors(hipStreamCreate(&(streams[i]))); + checkCudaErrors(cudaStreamCreate(&(streams[i]))); + } + + // create CUDA event handles + cudaEvent_t start_event, stop_event; + + // CHECK: checkCudaErrors(hipEventCreate(&start_event)); + // CHECK: checkCudaErrors(hipEventCreate(&stop_event)); + checkCudaErrors(cudaEventCreate(&start_event)); + checkCudaErrors(cudaEventCreate(&stop_event)); + + // the events are used for synchronization only and hence do not need to record timings + // this also makes events not introduce global sync points when recorded which is critical to get overlap + + // CHECK: hipEvent_t *kernelEvent; + // CHECK: kernelEvent = (hipEvent_t *) malloc(nkernels * sizeof(hipEvent_t)); + cudaEvent_t *kernelEvent; + kernelEvent = (cudaEvent_t *) malloc(nkernels * sizeof(cudaEvent_t)); + + for (int i = 0; i < nkernels; i++) + { + // CHECK: checkCudaErrors(hipEventCreateWithFlags(&(kernelEvent[i]), hipEventDisableTiming)); + checkCudaErrors(cudaEventCreateWithFlags(&(kernelEvent[i]), cudaEventDisableTiming)); + } + + ////////////////////////////////////////////////////////////////////// + // time execution with nkernels streams + clock_t total_clocks = 0; +#if defined(__arm__) || defined(__aarch64__) + // the kernel takes more time than the channel reset time on arm archs, so to prevent hangs reduce time_clocks. + clock_t time_clocks = (clock_t)(kernel_time * (deviceProp.clockRate / 1000)); +#else + clock_t time_clocks = (clock_t)(kernel_time * deviceProp.clockRate); +#endif + + // CHECK: hipEventRecord(start_event, 0); + cudaEventRecord(start_event, 0); + + // queue nkernels in separate streams and record when they are done + for (int i=0; i>>(&d_a[i], time_clocks); + total_clocks += time_clocks; + + // CHECK: checkCudaErrors(hipEventRecord(kernelEvent[i], streams[i])); + checkCudaErrors(cudaEventRecord(kernelEvent[i], streams[i])); + + // make the last stream wait for the kernel event to be recorded + // CHECK: checkCudaErrors(hipStreamWaitEvent(streams[nstreams-1], kernelEvent[i],0)); + checkCudaErrors(cudaStreamWaitEvent(streams[nstreams-1], kernelEvent[i],0)); + } + + // queue a sum kernel and a copy back to host in the last stream. + // the commands in this stream get dispatched as soon as all the kernel events have been recorded + // CHECK: hipLaunchKernelGGL(sum, dim3(1), dim3(32), 0, streams[nstreams-1], d_a, nkernels); + // CHECK: checkCudaErrors(hipMemcpyAsync(a, d_a, sizeof(clock_t), hipMemcpyDeviceToHost, streams[nstreams-1])); + sum<<<1,32,0,streams[nstreams-1]>>>(d_a, nkernels); + checkCudaErrors(cudaMemcpyAsync(a, d_a, sizeof(clock_t), cudaMemcpyDeviceToHost, streams[nstreams-1])); + + // at this point the CPU has dispatched all work for the GPU and can continue processing other tasks in parallel + + // in this sample we just wait until the GPU is done + // CHECK: checkCudaErrors(hipEventRecord(stop_event, 0)); + // CHECK: checkCudaErrors(hipEventSynchronize(stop_event)); + // CHECK: checkCudaErrors(hipEventElapsedTime(&elapsed_time, start_event, stop_event)); + checkCudaErrors(cudaEventRecord(stop_event, 0)); + checkCudaErrors(cudaEventSynchronize(stop_event)); + checkCudaErrors(cudaEventElapsedTime(&elapsed_time, start_event, stop_event)); + + printf("Expected time for serial execution of %d kernels = %.3fs\n", nkernels, nkernels * kernel_time/1000.0f); + printf("Expected time for concurrent execution of %d kernels = %.3fs\n", nkernels, kernel_time/1000.0f); + printf("Measured time for sample = %.3fs\n", elapsed_time/1000.0f); + + bool bTestResult = (a[0] > total_clocks); + + // release resources + for (int i = 0; i < nkernels; i++) + { + // CHECK: hipStreamDestroy(streams[i]); + // CHECK: hipEventDestroy(kernelEvent[i]); + cudaStreamDestroy(streams[i]); + cudaEventDestroy(kernelEvent[i]); + } + + free(streams); + free(kernelEvent); + + // CHECK: hipEventDestroy(start_event); + // CHECK: hipEventDestroy(stop_event); + // CHECK: hipHostFree(a); + // CHECK: hipFree(d_a); + cudaEventDestroy(start_event); + cudaEventDestroy(stop_event); + cudaFreeHost(a); + cudaFree(d_a); + + if (!bTestResult) + { + printf("Test failed!\n"); + exit(EXIT_FAILURE); + } + + printf("Test passed\n"); + exit(EXIT_SUCCESS); +} From cdd849541f6a5f17c6db4c16b8d432b82bd1033e Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 23 Oct 2017 21:16:13 +0300 Subject: [PATCH 28/32] [HIPIFY] cmake: fix standalone build --- hipamd/hipify-clang/CMakeLists.txt | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) diff --git a/hipamd/hipify-clang/CMakeLists.txt b/hipamd/hipify-clang/CMakeLists.txt index 355b0cea7a..2d9cbb1feb 100644 --- a/hipamd/hipify-clang/CMakeLists.txt +++ b/hipamd/hipify-clang/CMakeLists.txt @@ -75,15 +75,21 @@ if (HIPIFY_CLANG_TESTS) set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) + if (CMAKE_CURRENT_SOURCE_DIR EQUAL CMAKE_SOURCE_DIR) + set(SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}) + else() + set(SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/..) + endif() + configure_file( - ${CMAKE_SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in - ${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg + ${SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in + ${BINARY_DIR}/tests/hipify-clang/lit.site.cfg @ONLY ) add_lit_testsuite(test-hipify "Running HIPify regression tests" - ${CMAKE_SOURCE_DIR}/tests/hipify-clang - PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg + ${SOURCE_DIR}/tests/hipify-clang + PARAMS site_config=${BINARY_DIR}/tests/hipify-clang/lit.site.cfg DEPENDS hipify-clang lit ) From aa288fcf1fab4b7426ef88e3730cfbe65d63c6cc Mon Sep 17 00:00:00 2001 From: Chris Kitching Date: Tue, 24 Oct 2017 01:45:23 +0100 Subject: [PATCH 29/32] Tweak some version numbers in clang version compatibility checks Apparently a couple of those APIs changed in clang 5, not 4. Drat. --- hipamd/hipify-clang/src/Cuda2Hip.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hipamd/hipify-clang/src/Cuda2Hip.cpp b/hipamd/hipify-clang/src/Cuda2Hip.cpp index 0f73859893..1f0bcdba65 100644 --- a/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -3031,7 +3031,7 @@ public: : Cuda2Hip(R, mainFileName) {} virtual bool handleBeginSource(CompilerInstance &CI -#if LLVM_VERSION_MAJOR < 4 +#if LLVM_VERSION_MAJOR <= 4 , StringRef Filename #endif ) override { @@ -3114,7 +3114,7 @@ public: if (_sm->isWrittenInMainFile(MacroNameTok.getLocation())) { // The getNumArgs function was rather unhelpfully renamed in clang 4.0. Its semantics // remain unchanged. -#if LLVM_VERSION_MAJOR > 3 +#if LLVM_VERSION_MAJOR > 4 #define GET_NUM_ARGS() getNumParams() #else #define GET_NUM_ARGS() getNumArgs() From 559d9a68aab1a40c3ca60d22303b3d84b2491a0e Mon Sep 17 00:00:00 2001 From: emankov Date: Tue, 24 Oct 2017 10:51:11 +0300 Subject: [PATCH 30/32] [HIPIFY] cmake: simplify build --- hipamd/hipify-clang/CMakeLists.txt | 16 ++++------------ 1 file changed, 4 insertions(+), 12 deletions(-) diff --git a/hipamd/hipify-clang/CMakeLists.txt b/hipamd/hipify-clang/CMakeLists.txt index 2d9cbb1feb..5f4f96fb8a 100644 --- a/hipamd/hipify-clang/CMakeLists.txt +++ b/hipamd/hipify-clang/CMakeLists.txt @@ -73,23 +73,15 @@ if (HIPIFY_CLANG_TESTS) # tests find_package(PythonInterp 2.7 REQUIRED EXACT) - set(BINARY_DIR ${CMAKE_CURRENT_BINARY_DIR}) - - if (CMAKE_CURRENT_SOURCE_DIR EQUAL CMAKE_SOURCE_DIR) - set(SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}) - else() - set(SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR}/..) - endif() - configure_file( - ${SOURCE_DIR}/tests/hipify-clang/lit.site.cfg.in - ${BINARY_DIR}/tests/hipify-clang/lit.site.cfg + $${CMAKE_CURRENT_LIST_DIR}/../tests/hipify-clang/lit.site.cfg.in + ${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg @ONLY ) add_lit_testsuite(test-hipify "Running HIPify regression tests" - ${SOURCE_DIR}/tests/hipify-clang - PARAMS site_config=${BINARY_DIR}/tests/hipify-clang/lit.site.cfg + ${CMAKE_CURRENT_LIST_DIR}/../tests/hipify-clang + PARAMS site_config=${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg DEPENDS hipify-clang lit ) From 7c51376bbf0238885d21f07e6adab40d10c3a515 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 24 Oct 2017 14:16:05 +0300 Subject: [PATCH 31/32] [HIPIFY][fix] cmake: do not build hipify-clang if not asked + warn "hipify-clang will not be built" if HIPIFY_CLANG_LLVM_DIR is not specified. + fix typo in previous commit . --- hipamd/hipify-clang/CMakeLists.txt | 15 ++++++++++----- 1 file changed, 10 insertions(+), 5 deletions(-) diff --git a/hipamd/hipify-clang/CMakeLists.txt b/hipamd/hipify-clang/CMakeLists.txt index 5f4f96fb8a..b9d009ad92 100644 --- a/hipamd/hipify-clang/CMakeLists.txt +++ b/hipamd/hipify-clang/CMakeLists.txt @@ -1,6 +1,15 @@ cmake_minimum_required(VERSION 2.8.8) project(hipify-clang) +set(BUILD_HIPIFY_CLANG 0 CACHE INTERNAL "") + +if (HIPIFY_CLANG_LLVM_DIR) + find_package(LLVM PATHS ${HIPIFY_CLANG_LLVM_DIR} REQUIRED) +else() + message(STATUS "hipify-clang will not be built. To build it please specify absolute path to LLVM 3.8 or higher using HIPIFY_CLANG_LLVM_DIR") + return() +endif() + option(HIPIFY_CLANG_TESTS "Build the tests for hipify-clang, if lit is installed" ON) # Disable the tests if `lit` is not installed. @@ -10,10 +19,6 @@ if (NOT LIT_COMMAND) message(STATUS "hipify-clang's tests are not being built because `lit` is not installed.") endif() -set(BUILD_HIPIFY_CLANG 0 CACHE INTERNAL "") - -find_package(LLVM PATHS ${HIPIFY_CLANG_LLVM_DIR} REQUIRED) - list(APPEND CMAKE_MODULE_PATH ${LLVM_CMAKE_DIR}) include(AddLLVM) @@ -74,7 +79,7 @@ if (HIPIFY_CLANG_TESTS) find_package(PythonInterp 2.7 REQUIRED EXACT) configure_file( - $${CMAKE_CURRENT_LIST_DIR}/../tests/hipify-clang/lit.site.cfg.in + ${CMAKE_CURRENT_LIST_DIR}/../tests/hipify-clang/lit.site.cfg.in ${CMAKE_CURRENT_BINARY_DIR}/tests/hipify-clang/lit.site.cfg @ONLY ) From 4a0228cedbcdd41c37e06c3015798c0daefeafe2 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 24 Oct 2017 16:35:10 +0300 Subject: [PATCH 32/32] [HIPIFY][fix] cmake: NO_DEFAULT_PATH is strongly needed in find_package for LLVM Otherwise LLVM will be searched in system folders. --- hipamd/hipify-clang/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/hipify-clang/CMakeLists.txt b/hipamd/hipify-clang/CMakeLists.txt index b9d009ad92..7bc2ee9367 100644 --- a/hipamd/hipify-clang/CMakeLists.txt +++ b/hipamd/hipify-clang/CMakeLists.txt @@ -4,7 +4,7 @@ project(hipify-clang) set(BUILD_HIPIFY_CLANG 0 CACHE INTERNAL "") if (HIPIFY_CLANG_LLVM_DIR) - find_package(LLVM PATHS ${HIPIFY_CLANG_LLVM_DIR} REQUIRED) + find_package(LLVM PATHS ${HIPIFY_CLANG_LLVM_DIR} REQUIRED NO_DEFAULT_PATH) else() message(STATUS "hipify-clang will not be built. To build it please specify absolute path to LLVM 3.8 or higher using HIPIFY_CLANG_LLVM_DIR") return()