From ec03a8b3520235b085bfc915f1047846ad6a7f91 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Tue, 18 Dec 2018 15:48:21 -0500 Subject: [PATCH 01/18] Add default arguments for hipConfigureCall --- include/hip/hcc_detail/hip_runtime_api.h | 2 +- tests/src/kernel/hipGridLaunch.cpp | 31 ++++++++++++++++++++++++ 2 files changed, 32 insertions(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 97d6320f7f..d4b8c1faf4 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2573,7 +2573,7 @@ hipError_t hipIpcCloseMemHandle(void* devPtr); * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue * */ -hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem, hipStream_t stream); +hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem = 0, hipStream_t stream = 0); /** diff --git a/tests/src/kernel/hipGridLaunch.cpp b/tests/src/kernel/hipGridLaunch.cpp index 8505a0758f..4f7883c656 100644 --- a/tests/src/kernel/hipGridLaunch.cpp +++ b/tests/src/kernel/hipGridLaunch.cpp @@ -74,12 +74,43 @@ int test_gl2(size_t N) { return 0; } +#if __HIP__ +int test_triple_chevron(size_t N) { + size_t Nbytes = N * sizeof(int); + + int *A_d, *B_d, *C_d; + int *A_h, *B_h, *C_h; + + HipTest::initArrays(&A_d, &B_d, &C_d, &A_h, &B_h, &C_h, N); + + + unsigned blocks = HipTest::setNumBlocks(blocksPerCU, threadsPerBlock, N); + + + // Full vadd in one large chunk, to get things started: + HIPCHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice)); + HIPCHECK(hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice)); + + vectorADD2<<>>(A_d, B_d, C_d, N); + + HIPCHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); + + HIPCHECK(hipDeviceSynchronize()); + + HipTest::checkVectorADD(A_h, B_h, C_h, N); + + return 0; +} +#endif int main(int argc, char* argv[]) { HipTest::parseStandardArguments(argc, argv, true); test_gl2(N); +#if __HIP__ + test_triple_chevron(N); +#endif passed(); } From 4272119eeb8e8d277e4b44f2c893de4dda8c350b Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Thu, 20 Dec 2018 06:51:22 +0530 Subject: [PATCH 02/18] Remove redundant arguments from hipMemset3D test --- tests/src/runtimeApi/memory/hipMemset3D.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/tests/src/runtimeApi/memory/hipMemset3D.cpp b/tests/src/runtimeApi/memory/hipMemset3D.cpp index 53e21c93bb..151a86a6c0 100644 --- a/tests/src/runtimeApi/memory/hipMemset3D.cpp +++ b/tests/src/runtimeApi/memory/hipMemset3D.cpp @@ -24,8 +24,7 @@ THE SOFTWARE. /* HIT_START * BUILD: %t %s ../../test_common.cpp - * //Small copy - * RUN: %t -N 10 --memsetval 0x42 + * RUN: %t * HIT_END */ From 208d7cfaf14c433cbc9a253db90db82f60591c89 Mon Sep 17 00:00:00 2001 From: emankov Date: Fri, 28 Dec 2018 01:34:35 +0300 Subject: [PATCH 03/18] [HIPIFY] LLVM compatibility + sys::fs::make_absolute fro LLVM < 5.0 + sys::fs::real_path for LLVM >= 5.0 --- hipify-clang/src/LLVMCompat.cpp | 9 +++++++++ hipify-clang/src/LLVMCompat.h | 3 +++ hipify-clang/src/main.cpp | 6 +++--- 3 files changed, 15 insertions(+), 3 deletions(-) diff --git a/hipify-clang/src/LLVMCompat.cpp b/hipify-clang/src/LLVMCompat.cpp index f96bd86bb0..22a0b56d64 100644 --- a/hipify-clang/src/LLVMCompat.cpp +++ b/hipify-clang/src/LLVMCompat.cpp @@ -94,4 +94,13 @@ clang::SourceLocation getEndLoc(const clang::TypeLoc& typeLoc) { #endif } +std::error_code real_path(const Twine &path, SmallVectorImpl &output, + bool expand_tilde) { +#if LLVM_VERSION_MAJOR < 5 + return sys::fs::make_absolute(path, output); +#else + return sys::fs::real_path(path, output, expand_tilde); +#endif +} + } // namespace llcompat diff --git a/hipify-clang/src/LLVMCompat.h b/hipify-clang/src/LLVMCompat.h index 4d6edcc9f8..98ff66fb41 100644 --- a/hipify-clang/src/LLVMCompat.h +++ b/hipify-clang/src/LLVMCompat.h @@ -78,4 +78,7 @@ void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, size_t len, bool DisableMacroExpansion); +std::error_code real_path(const Twine &path, SmallVectorImpl &output, + bool expand_tilde = false); + } // namespace llcompat diff --git a/hipify-clang/src/main.cpp b/hipify-clang/src/main.cpp index 06981a9d70..7dc42f8882 100644 --- a/hipify-clang/src/main.cpp +++ b/hipify-clang/src/main.cpp @@ -46,7 +46,7 @@ std::string getAbsoluteDirectory(const std::string& sDir, std::error_code& EC, return sDir; } SmallString<256> dirAbsPath; - EC = sys::fs::real_path(sDir, dirAbsPath, true); + EC = llcompat::real_path(sDir, dirAbsPath, true); if (!EC && sys::fs::is_regular_file(dirAbsPath)) { llvm::errs() << "\n" << sHipify << sError << sDir << " is not a directory\n"; EC = std::error_code(static_cast(std::errc::not_a_directory), std::generic_category()); @@ -58,7 +58,7 @@ std::string getAbsoluteDirectory(const std::string& sDir, std::error_code& EC, llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << sDirType << " directory: " << sDir << "\n"; return ""; } - EC = sys::fs::real_path(sDir, dirAbsPath, true); + EC = llcompat::real_path(sDir, dirAbsPath, true); if (EC) { llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << sDirType << " directory: " << sDir << "\n"; return ""; @@ -137,7 +137,7 @@ int main(int argc, const char **argv) { // Create a copy of the file to work on. When we're done, we'll move this onto the // output (which may mean overwriting the input, if we're in-place). // Should we fail for some reason, we'll just leak this file and not corrupt the input. - EC = sys::fs::real_path(src, sourceAbsPath, true); + EC = llcompat::real_path(src, sourceAbsPath, true); if (EC) { llvm::errs() << "\n" << sHipify << sError << EC.message() << ": " << src << "\n"; Result = 1; From d5d6541ad5a9fd5dfa7b355c87eedd3164414d70 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 28 Dec 2018 11:23:52 +0530 Subject: [PATCH 04/18] [ci] Remove ROCm 1.9.x from test infrastructure Change-Id: I58cce7c60f6def22f337bbe15fab2620419e202b --- Jenkinsfile | 47 +---------------------------------------------- 1 file changed, 1 insertion(+), 46 deletions(-) diff --git a/Jenkinsfile b/Jenkinsfile index 1257bfb0b4..812b68633a 100644 --- a/Jenkinsfile +++ b/Jenkinsfile @@ -295,52 +295,7 @@ String build_config = 'Release' String job_name = env.JOB_NAME.toLowerCase( ) // The following launches 3 builds in parallel: rocm-head, rocm-1.9.x and cuda-9.x -parallel rocm_1_9: -{ - node('hip-rocm') - { - String hcc_ver = 'rocm-1.9.x' - String from_image = 'ci_test_nodes/rocm-1.9.x/ubuntu-16.04:latest' - String inside_args = '--device=/dev/kfd --device=/dev/dri --group-add=video' - - // Checkout source code, dependencies and version files - String source_hip_rel = checkout_and_version( hcc_ver ) - - // Create/reuse a docker image that represents the hip build environment - def hip_build_image = docker_build_image( hcc_ver, 'hip', '', source_hip_rel, from_image ) - - // Print system information for the log - hip_build_image.inside( inside_args ) - { - sh """#!/usr/bin/env bash - set -x - /opt/rocm/bin/rocm_agent_enumerator -t ALL - /opt/rocm/bin/hcc --version - """ - } - - // Conctruct a binary directory path based on build config - String build_hip_rel = build_directory_rel( build_config ); - - // Build hip inside of the build environment - docker_build_inside_image( hip_build_image, inside_args, hcc_ver, '', build_config, source_hip_rel, build_hip_rel ) - - // Clean docker build image - docker_clean_images( 'hip', docker_build_image_name( ) ) - - // After a successful build, upload a docker image of the results - /* - String hip_image_name = docker_upload_artifactory( hcc_ver, job_name, from_image, source_hip_rel, build_hip_rel ) - if( params.push_image_to_docker_hub ) - { - docker_upload_dockerhub( job_name, hip_image_name, 'rocm' ) - docker_clean_images( 'rocm', hip_image_name ) - } - docker_clean_images( job_name, hip_image_name ) - */ - } -}, -rocm_2_0: +parallel rocm_2_0: { node('hip-rocm') { From 2d3f62c3e11668030a026f243982f3728c6c7b40 Mon Sep 17 00:00:00 2001 From: emankov Date: Sat, 29 Dec 2018 15:51:06 +0300 Subject: [PATCH 05/18] [HIPIFY][win] Set -std=c++14 if MSVC + Due to latest Windows SDK implementation based on c++14 features, '-std+c++14' is a must; + Doesn't affect older versions of MSVC 2017, 2015 (the latest supported is MSVC 2015); + On linux still '-std=c++11' is a must; + Exclude '-std=' from lit config. --- hipify-clang/src/main.cpp | 6 +++++- tests/hipify-clang/lit.cfg | 2 +- 2 files changed, 6 insertions(+), 2 deletions(-) diff --git a/hipify-clang/src/main.cpp b/hipify-clang/src/main.cpp index 7dc42f8882..8549e80d18 100644 --- a/hipify-clang/src/main.cpp +++ b/hipify-clang/src/main.cpp @@ -173,7 +173,11 @@ int main(int argc, const char **argv) { Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sInclude.c_str(), ct::ArgumentInsertPosition::BEGIN)); Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN)); // Ensure at least c++11 is used. - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-std=c++11", ct::ArgumentInsertPosition::BEGIN)); + std::string stdCpp = "-std=c++11"; +#if defined(_MSC_VER) + stdCpp = "-std=c++14"; +#endif + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(stdCpp.c_str(), ct::ArgumentInsertPosition::BEGIN)); #if defined(HIPIFY_CLANG_RES) Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES)); #endif diff --git a/tests/hipify-clang/lit.cfg b/tests/hipify-clang/lit.cfg index d0328c06c2..e5805a5663 100644 --- a/tests/hipify-clang/lit.cfg +++ b/tests/hipify-clang/lit.cfg @@ -65,7 +65,7 @@ clang_args = "-x cuda -v --cuda-gpu-arch=sm_30 --cuda-path='%s'" if sys.platform in ['win32']: run_test_ext = ".bat" hipify_path += "/" + config.build_type - clang_args += " -isystem'%s'/common/inc -std=c++14" + clang_args += " -isystem'%s'/common/inc" else: run_test_ext = ".sh" clang_args += " -isystem'%s'/samples/common/inc" From 7ea586c323ee0c3df0f6d9d09937e215a31c023e Mon Sep 17 00:00:00 2001 From: emankov Date: Sat, 29 Dec 2018 17:04:59 +0300 Subject: [PATCH 06/18] [HIPIFY] Starts implicitly setting '-x cuda' by hipify-clang itself + No need in setting '-x cuda' for the user anymore; + Testing and Readme updated accordingly. --- hipify-clang/README.md | 1 - hipify-clang/src/main.cpp | 2 ++ tests/hipify-clang/lit.cfg | 2 +- 3 files changed, 3 insertions(+), 2 deletions(-) diff --git a/hipify-clang/README.md b/hipify-clang/README.md index 2df3a73bff..845d0e0a9e 100644 --- a/hipify-clang/README.md +++ b/hipify-clang/README.md @@ -321,7 +321,6 @@ For example: ./hipify-clang \ square.cu \ -- \ - -x cuda \ --cuda-path=/usr/local/cuda-8.0 \ --cuda-gpu-arch=sm_50 \ -isystem /usr/local/cuda-8.0/samples/common/inc diff --git a/hipify-clang/src/main.cpp b/hipify-clang/src/main.cpp index 8549e80d18..a8f7e6a2f9 100644 --- a/hipify-clang/src/main.cpp +++ b/hipify-clang/src/main.cpp @@ -171,6 +171,8 @@ int main(int argc, const char **argv) { ReplacementsFrontendActionFactory actionFactory(&replacementsToUse); std::string sInclude = "-I" + sys::path::parent_path(sourceAbsPath).str(); Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster(sInclude.c_str(), ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("cuda", ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-x", ct::ArgumentInsertPosition::BEGIN)); Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN)); // Ensure at least c++11 is used. std::string stdCpp = "-std=c++11"; diff --git a/tests/hipify-clang/lit.cfg b/tests/hipify-clang/lit.cfg index e5805a5663..db1508a38f 100644 --- a/tests/hipify-clang/lit.cfg +++ b/tests/hipify-clang/lit.cfg @@ -60,7 +60,7 @@ if obj_root is not None: config.environment['PATH'] = path hipify_path = obj_root -clang_args = "-x cuda -v --cuda-gpu-arch=sm_30 --cuda-path='%s'" +clang_args = "-v --cuda-gpu-arch=sm_30 --cuda-path='%s'" if sys.platform in ['win32']: run_test_ext = ".bat" From dec459efcadb204161e00fdb9e185cc349364263 Mon Sep 17 00:00:00 2001 From: emankov Date: Sat, 29 Dec 2018 19:48:28 +0300 Subject: [PATCH 07/18] [HIPIFY] Get rid of setting '--cuda-gpu-arch=' [Reasons] + We don't compile kernel code at least for now as HIP kernel syntax is almost equal CUDA's; + clang always includes PTX in its binaries, so e.g. a binary compiled with --cuda-gpu-arch= would be forwards-compatible with e.g. sm_35 GPUs. --- hipify-clang/README.md | 3 +-- tests/hipify-clang/lit.cfg | 2 +- 2 files changed, 2 insertions(+), 3 deletions(-) diff --git a/hipify-clang/README.md b/hipify-clang/README.md index 845d0e0a9e..5b4424b791 100644 --- a/hipify-clang/README.md +++ b/hipify-clang/README.md @@ -322,7 +322,6 @@ For example: square.cu \ -- \ --cuda-path=/usr/local/cuda-8.0 \ - --cuda-gpu-arch=sm_50 \ -isystem /usr/local/cuda-8.0/samples/common/inc ``` @@ -338,5 +337,5 @@ The information contained herein is for informational purposes only, and is subj AMD, the AMD Arrow logo, and combinations thereof are trademarks of Advanced Micro Devices, Inc. Other product names used in this publication are for identification purposes only and may be trademarks of their respective companies. -Copyright (c) 2014-2018 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2014-2019 Advanced Micro Devices, Inc. All rights reserved. diff --git a/tests/hipify-clang/lit.cfg b/tests/hipify-clang/lit.cfg index db1508a38f..7b25731ba6 100644 --- a/tests/hipify-clang/lit.cfg +++ b/tests/hipify-clang/lit.cfg @@ -60,7 +60,7 @@ if obj_root is not None: config.environment['PATH'] = path hipify_path = obj_root -clang_args = "-v --cuda-gpu-arch=sm_30 --cuda-path='%s'" +clang_args = "-v --cuda-path='%s'" if sys.platform in ['win32']: run_test_ext = ".bat" From c519c89abffd5eaa50ee5ff4e29b004fc9782e6c Mon Sep 17 00:00:00 2001 From: emankov Date: Sat, 29 Dec 2018 21:53:18 +0300 Subject: [PATCH 08/18] [HIPIFY] Add system includes for packaged hipify-clang to use it without the necessity of installing clang Header files in those dirs are an integral part of clang, which includes wrapper headers for CUDA. There is no need in checking those dirs. [Reasons] 1. clang will not take into account empty or irrelevant include dirs; 2. Packages for different clang versions will have different header files in the above dirs. --- hipify-clang/src/main.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/hipify-clang/src/main.cpp b/hipify-clang/src/main.cpp index a8f7e6a2f9..d2d3d074b4 100644 --- a/hipify-clang/src/main.cpp +++ b/hipify-clang/src/main.cpp @@ -174,6 +174,11 @@ int main(int argc, const char **argv) { Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("cuda", ct::ArgumentInsertPosition::BEGIN)); Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-x", ct::ArgumentInsertPosition::BEGIN)); Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN)); + // Includes for clang's CUDA wrappers for using by packaged hipify-clang + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include/cuda_wrappers", ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-isystem", ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include", ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-isystem", ct::ArgumentInsertPosition::BEGIN)); // Ensure at least c++11 is used. std::string stdCpp = "-std=c++11"; #if defined(_MSC_VER) From c09022232a7732bd70b4f39dd4adce8d539b3aa5 Mon Sep 17 00:00:00 2001 From: emankov Date: Sat, 29 Dec 2018 22:10:00 +0300 Subject: [PATCH 09/18] [HIPIFY][fix] Change system includes order 'include/cuda_wrappers' should go first. --- hipify-clang/src/main.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hipify-clang/src/main.cpp b/hipify-clang/src/main.cpp index d2d3d074b4..2a3b86f204 100644 --- a/hipify-clang/src/main.cpp +++ b/hipify-clang/src/main.cpp @@ -175,10 +175,10 @@ int main(int argc, const char **argv) { Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-x", ct::ArgumentInsertPosition::BEGIN)); Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN)); // Includes for clang's CUDA wrappers for using by packaged hipify-clang - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include/cuda_wrappers", ct::ArgumentInsertPosition::BEGIN)); - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-isystem", ct::ArgumentInsertPosition::BEGIN)); Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include", ct::ArgumentInsertPosition::BEGIN)); Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-isystem", ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("./include/cuda_wrappers", ct::ArgumentInsertPosition::BEGIN)); + Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-isystem", ct::ArgumentInsertPosition::BEGIN)); // Ensure at least c++11 is used. std::string stdCpp = "-std=c++11"; #if defined(_MSC_VER) From 548f4dd4e51250a579d1b5b57154ea2368be1fca Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Mon, 31 Dec 2018 13:07:19 -0500 Subject: [PATCH 10/18] Use __dparm for default parameter --- include/hip/hcc_detail/hip_runtime_api.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index d4b8c1faf4..0d526ca25d 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2573,7 +2573,7 @@ hipError_t hipIpcCloseMemHandle(void* devPtr); * @returns hipSuccess, hipInvalidDevice, hipErrorNotInitialized, hipErrorInvalidValue * */ -hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem = 0, hipStream_t stream = 0); +hipError_t hipConfigureCall(dim3 gridDim, dim3 blockDim, size_t sharedMem __dparm(0), hipStream_t stream __dparm(0)); /** From 25c7e5d6096330075559a1fea50a58ffb1a41295 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Tue, 18 Dec 2018 23:05:39 +0000 Subject: [PATCH 11/18] Start re-working 731 for 2.0. --- .../hip/hcc_detail/functional_grid_launch.hpp | 38 +++++-- include/hip/hcc_detail/program_state.hpp | 2 + src/program_state.cpp | 104 ++++++++++++++++++ 3 files changed, 137 insertions(+), 7 deletions(-) diff --git a/include/hip/hcc_detail/functional_grid_launch.hpp b/include/hip/hcc_detail/functional_grid_launch.hpp index e678f25aa2..cd90aa401a 100644 --- a/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/include/hip/hcc_detail/functional_grid_launch.hpp @@ -33,6 +33,7 @@ THE SOFTWARE. #include #include +#include #include #include #include @@ -56,7 +57,9 @@ template < typename... Ts, typename std::enable_if::type* = nullptr> inline std::vector make_kernarg( - std::vector kernarg, const std::tuple&) { + const std::tuple&, + const std::vector>&, + std::vector kernarg) { return kernarg; } @@ -65,7 +68,9 @@ template < typename... Ts, typename std::enable_if::type* = nullptr> inline std::vector make_kernarg( - std::vector kernarg, const std::tuple& formals) { + const std::tuple& formals, + const std::vector>& size_align, + std::vector kernarg) { using T = typename std::tuple_element>::type; static_assert( @@ -80,24 +85,43 @@ inline std::vector make_kernarg( #endif kernarg.resize(round_up_to_next_multiple_nonnegative( - kernarg.size(), alignof(T)) + sizeof(T)); + kernarg.size(), size_align[n].second) + size_align[n].first); - new (kernarg.data() + kernarg.size() - sizeof(T)) T{std::get(formals)}; + std::memcpy( + kernarg.data() + kernarg.size() - size_align[n].first, + &std::get(formals), + size_align[n].first); - return make_kernarg(std::move(kernarg), formals); + return make_kernarg(formals, size_align, std::move(kernarg)); } template inline std::vector make_kernarg( - void (*)(Formals...), std::tuple actuals) { + void (*kernel)(Formals...), std::tuple actuals) { static_assert(sizeof...(Formals) == sizeof...(Actuals), "The count of formal arguments must match the count of actuals."); + if (sizeof...(Formals) == 0) return {}; + + const auto it = function_names().find( + reinterpret_cast(kernel)); + + if (it == function_names().cend()) { + throw std::runtime_error{"Undefined __global__ function."}; + } + + const auto it1 = kernargs().find(it->second); + + if (it1 == kernargs().end()) { + throw std::runtime_error{ + "Missing metadata for __global__ function: " + it->second}; + } + std::tuple to_formals{std::move(actuals)}; std::vector kernarg; kernarg.reserve(sizeof(to_formals)); - return make_kernarg<0>(std::move(kernarg), to_formals); + return make_kernarg<0>(to_formals, it1->second, std::move(kernarg)); } void hipLaunchKernelGGLImpl(std::uintptr_t function_address, const dim3& numBlocks, diff --git a/include/hip/hcc_detail/program_state.hpp b/include/hip/hcc_detail/program_state.hpp index bdb87b3509..92bef22172 100644 --- a/include/hip/hcc_detail/program_state.hpp +++ b/include/hip/hcc_detail/program_state.hpp @@ -99,6 +99,8 @@ const std::unordered_map& function_names(bool rebuild = false); std::unordered_map& globals(bool rebuild = false); +std::unordered_map< + std::string, std::vector>>& kernargs(); hsa_executable_t load_executable(const std::string& file, hsa_executable_t executable, hsa_agent_t agent); diff --git a/src/program_state.cpp b/src/program_state.cpp index 88cdeeb404..922d827be4 100644 --- a/src/program_state.cpp +++ b/src/program_state.cpp @@ -340,6 +340,90 @@ void load_code_object_and_freeze_executable( code_readers.push_back(move(tmp)); } } + +size_t parse_args( + const string& metadata, + size_t f, + size_t l, + vector>& size_align) { + if (f == l) return f; + if (!size_align.empty()) return l; + + do { + static constexpr size_t size_sz{5}; + f = metadata.find("Size:", f) + size_sz; + + if (l <= f) return f; + + auto size = strtoul(&metadata[f], nullptr, 10); + + static constexpr size_t align_sz{6}; + f = metadata.find("Align:", f) + align_sz; + + char* l{}; + auto align = strtoul(&metadata[f], &l, 10); + + f += (l - &metadata[f]) + 1; + + size_align.emplace_back(size, align); + } while (true); +} + +void read_kernarg_metadata( + elfio& reader, + unordered_map>>& kernargs) +{ // TODO: this is inefficient. + auto it = find_section_if( + reader, [](const section* x) { return x->get_type() == SHT_NOTE; }); + + if (!it) return; + + const note_section_accessor acc{reader, it}; + for (decltype(acc.get_notes_num()) i = 0; i != acc.get_notes_num(); ++i) { + ELFIO::Elf_Word type{}; + string name{}; + void* desc{}; + Elf_Word desc_size{}; + + acc.get_note(i, type, name, desc, desc_size); + + if (name != "AMD") continue; // TODO: switch to using NT_AMD_AMDGPU_HSA_METADATA. + + string tmp{ + static_cast(desc), static_cast(desc) + desc_size}; + + auto dx = tmp.find("Kernels:"); + + if (dx == string::npos) continue; + + static constexpr decltype(tmp.size()) kernels_sz{8}; + dx += kernels_sz; + + do { + dx = tmp.find("Name:", dx); + + if (dx == string::npos) break; + + static constexpr decltype(tmp.size()) name_sz{5}; + dx = tmp.find_first_not_of(" '", dx + name_sz); + + auto fn = tmp.substr(dx, tmp.find_first_of("'\n", dx) - dx); + dx += fn.size(); + + auto dx1 = tmp.find("CodeProps", dx); + dx = tmp.find("Args:", dx); + + if (dx1 < dx) { + dx = dx1; + continue; + } + if (dx == string::npos) break; + + static constexpr decltype(tmp.size()) args_sz{5}; + dx = parse_args(tmp, dx + args_sz, dx1, kernargs[fn]); + } while (true); + } +} } // namespace namespace hip_impl { @@ -501,6 +585,26 @@ unordered_map& globals(bool rebuild) { return r; } + +unordered_map>>& kernargs() { + static unordered_map>> r; + static once_flag f; + + call_once(f, []() { + for (auto&& blob : code_object_blobs()) { + stringstream tmp{std::string{ + blob.second.front().cbegin(), blob.second.front().cend()}}; + + elfio reader; + if (!reader.load(tmp)) continue; + + read_kernarg_metadata(reader, r); + } + }); + + return r; +} + hsa_executable_t load_executable(const string& file, hsa_executable_t executable, hsa_agent_t agent) { elfio reader; From 340674ceb6ee471e50ed957b51d6a27847176c0f Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Wed, 19 Dec 2018 03:13:57 +0000 Subject: [PATCH 12/18] More blobs, more problems. --- src/program_state.cpp | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/src/program_state.cpp b/src/program_state.cpp index 922d827be4..38cae74dcc 100644 --- a/src/program_state.cpp +++ b/src/program_state.cpp @@ -591,14 +591,15 @@ unordered_map>>& kernargs() { static once_flag f; call_once(f, []() { - for (auto&& blob : code_object_blobs()) { - stringstream tmp{std::string{ - blob.second.front().cbegin(), blob.second.front().cend()}}; + for (auto&& isa_blobs : code_object_blobs()) { + for (auto&& blob : isa_blobs.second) { + stringstream tmp{std::string{blob.cbegin(), blob.cend()}}; - elfio reader; - if (!reader.load(tmp)) continue; + elfio reader; + if (!reader.load(tmp)) continue; - read_kernarg_metadata(reader, r); + read_kernarg_metadata(reader, r); + } } }); From ec14daa7cee0ea5c1073dfc3c369c5ab1a50b37d Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 20 Dec 2018 00:26:42 +0000 Subject: [PATCH 13/18] Hook into the creaky lazy-reinit machinery. Try to minimise race-risk. --- .../hip/hcc_detail/functional_grid_launch.hpp | 21 +++++++++----- include/hip/hcc_detail/program_state.hpp | 5 ++-- src/program_state.cpp | 29 +++++++++++++++---- 3 files changed, 40 insertions(+), 15 deletions(-) diff --git a/include/hip/hcc_detail/functional_grid_launch.hpp b/include/hip/hcc_detail/functional_grid_launch.hpp index cd90aa401a..2fbda48629 100644 --- a/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/include/hip/hcc_detail/functional_grid_launch.hpp @@ -103,18 +103,23 @@ inline std::vector make_kernarg( if (sizeof...(Formals) == 0) return {}; - const auto it = function_names().find( - reinterpret_cast(kernel)); - + auto it = function_names().find(reinterpret_cast(kernel)); if (it == function_names().cend()) { - throw std::runtime_error{"Undefined __global__ function."}; + it = + function_names(true).find(reinterpret_cast(kernel)); + if (it == function_names().cend()) { + throw std::runtime_error{"Undefined __global__ function."}; + } } - const auto it1 = kernargs().find(it->second); - + auto it1 = kernargs().find(it->second); if (it1 == kernargs().end()) { - throw std::runtime_error{ - "Missing metadata for __global__ function: " + it->second}; + it1 = kernargs(true).find(it->second); + + if (it1 == kernargs().end()) { + throw std::runtime_error{ + "Missing metadata for __global__ function: " + it->second}; + } } std::tuple to_formals{std::move(actuals)}; diff --git a/include/hip/hcc_detail/program_state.hpp b/include/hip/hcc_detail/program_state.hpp index 92bef22172..da13c7c3db 100644 --- a/include/hip/hcc_detail/program_state.hpp +++ b/include/hip/hcc_detail/program_state.hpp @@ -99,8 +99,9 @@ const std::unordered_map& function_names(bool rebuild = false); std::unordered_map& globals(bool rebuild = false); -std::unordered_map< - std::string, std::vector>>& kernargs(); +const std::unordered_map< + std::string, std::vector>>& + kernargs(bool rebuild = false); hsa_executable_t load_executable(const std::string& file, hsa_executable_t executable, hsa_agent_t agent); diff --git a/src/program_state.cpp b/src/program_state.cpp index 38cae74dcc..bb906b0ad9 100644 --- a/src/program_state.cpp +++ b/src/program_state.cpp @@ -538,6 +538,7 @@ const unordered_map>>& fu // created previously function_names(rebuild); + kernargs(rebuild); kernels(rebuild); globals(rebuild); } @@ -585,12 +586,12 @@ unordered_map& globals(bool rebuild) { return r; } - -unordered_map>>& kernargs() { +const unordered_map>>& kernargs( + bool rebuild) { static unordered_map>> r; static once_flag f; - call_once(f, []() { + static const auto build_map = [](decltype(r)& x) { for (auto&& isa_blobs : code_object_blobs()) { for (auto&& blob : isa_blobs.second) { stringstream tmp{std::string{blob.cbegin(), blob.cend()}}; @@ -598,10 +599,28 @@ unordered_map>>& kernargs() { elfio reader; if (!reader.load(tmp)) continue; - read_kernarg_metadata(reader, r); + read_kernarg_metadata(reader, x); } } - }); + }; + call_once(f, []() { r.reserve(function_names().size()); build_map(r); }); + + if (rebuild) { + static mutex mtx; + thread_local static decltype(r) tmp; + + { + lock_guard lck{mtx}; + + tmp.insert(r.cbegin(), r.cend()); // Should use merge in C++17. + } + + build_map(tmp); + + lock_guard lck{mtx}; + + r.insert(tmp.cbegin(), tmp.cend()); + } return r; } From 7fb0f4ee2a4a5a2ce933b99796054035d595bf43 Mon Sep 17 00:00:00 2001 From: emankov Date: Thu, 3 Jan 2019 17:17:21 +0300 Subject: [PATCH 14/18] [HIPIFY][tests] Reduce cudaRegister.cu test --- tests/hipify-clang/unit_tests/samples/cudaRegister.cu | 8 -------- 1 file changed, 8 deletions(-) diff --git a/tests/hipify-clang/unit_tests/samples/cudaRegister.cu b/tests/hipify-clang/unit_tests/samples/cudaRegister.cu index 43b4345337..11192c452e 100644 --- a/tests/hipify-clang/unit_tests/samples/cudaRegister.cu +++ b/tests/hipify-clang/unit_tests/samples/cudaRegister.cu @@ -22,12 +22,6 @@ THE SOFTWARE. #include #include #include -#ifdef _WIN32 -#include -#define sleep(x) Sleep(x) -#else -#include -#endif #include #include @@ -90,7 +84,6 @@ int main(){ // CHECK: hipLaunchKernelGGL(Inc1, dim3(dimGrid), dim3(dimBlock), 0, 0, Ad, Bd); Inc1<<>>(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: Thu, 3 Jan 2019 18:06:49 +0300 Subject: [PATCH 15/18] [HIPIFY][tests] Exclude CUDA 8.0 (9.0) tests if CUDA < 8.0 (9.0) --- tests/hipify-clang/lit.cfg | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/tests/hipify-clang/lit.cfg b/tests/hipify-clang/lit.cfg index 7b25731ba6..462aba6c52 100644 --- a/tests/hipify-clang/lit.cfg +++ b/tests/hipify-clang/lit.cfg @@ -15,12 +15,14 @@ print("CUDA " + config.cuda_version + " will be used for testing.") config.excludes = ['cmdparser.hpp'] +if config.cuda_version_major < 8: + config.excludes.append('cuSPARSE_02.cu') if config.cuda_version_major < 9: config.excludes.append('cuSPARSE_04.cu') config.excludes.append('cuSPARSE_05.cu') config.excludes.append('cuSPARSE_06.cu') config.excludes.append('cuSPARSE_07.cu') - + config.excludes.append('benchmark_curand_kernel.cpp') if config.cuda_version_major < 10: config.excludes.append('cuSPARSE_08.cu') config.excludes.append('cuSPARSE_09.cu') From 9abcad9407d6c607a29516d349faf4d52ce7ed51 Mon Sep 17 00:00:00 2001 From: emankov Date: Fri, 4 Jan 2019 01:32:30 +0300 Subject: [PATCH 16/18] [HIPIFY][tests] CUDA 7.0 is also supported + Exclude 1 test for CUDA 7.0 + Update README.md --- hipify-clang/README.md | 2 +- tests/hipify-clang/lit.cfg | 2 ++ tests/hipify-clang/lit.site.cfg.in | 1 + 3 files changed, 4 insertions(+), 1 deletion(-) diff --git a/hipify-clang/README.md b/hipify-clang/README.md index 5b4424b791..e553fb4adc 100644 --- a/hipify-clang/README.md +++ b/hipify-clang/README.md @@ -34,7 +34,7 @@ `hipify-clang` requires: 1. LLVM+CLANG of at least version 3.8.0, latest stable and recommended release: 6.0.1 (linux and windows). -2. CUDA at least version 7.5, latest supported release is 9.0. +2. CUDA at least version 7.0, latest supported version is 9.0. | **LLVM release version** | **CUDA latest supported version** | **Comments** | |:------------------------:|:---------------------------------:|:------------:| diff --git a/tests/hipify-clang/lit.cfg b/tests/hipify-clang/lit.cfg index 462aba6c52..76a811031a 100644 --- a/tests/hipify-clang/lit.cfg +++ b/tests/hipify-clang/lit.cfg @@ -15,6 +15,8 @@ print("CUDA " + config.cuda_version + " will be used for testing.") config.excludes = ['cmdparser.hpp'] +if config.cuda_version_major == 7 and config.cuda_version_minor == 0: + config.excludes.append('headers_test_09.cu') if config.cuda_version_major < 8: config.excludes.append('cuSPARSE_02.cu') if config.cuda_version_major < 9: diff --git a/tests/hipify-clang/lit.site.cfg.in b/tests/hipify-clang/lit.site.cfg.in index f7302252a8..f93c206be4 100644 --- a/tests/hipify-clang/lit.site.cfg.in +++ b/tests/hipify-clang/lit.site.cfg.in @@ -6,6 +6,7 @@ config.obj_root = "@CMAKE_CURRENT_BINARY_DIR@" config.cuda_root = "@CUDA_TOOLKIT_ROOT_DIR@" config.cuda_dnn_root = "@CUDA_DNN_ROOT_DIR@" config.cuda_version_major = int("@CUDA_VERSION_MAJOR@") +config.cuda_version_minor = int("@CUDA_VERSION_MINOR@") config.cuda_version = "@CUDA_VERSION@" if sys.platform in ['win32']: config.cuda_sdk_root = "@CUDA_SDK_ROOT_DIR@" From cd15a856638d0a2dcc1c5c999d3baca37cdcafba Mon Sep 17 00:00:00 2001 From: emankov Date: Fri, 4 Jan 2019 18:52:28 +0300 Subject: [PATCH 17/18] [HIPIFY][cmake][win] check for Visual Studio version hipify-clang could be built by Visual Studio 14 2015 or higher. [Reason] hipify-clang contains c++14 features which are not supported even by the latest VS 2013 and the Platform Toolset "Visual C++ Compiler Nov 2013 CTP". --- hipify-clang/CMakeLists.txt | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/hipify-clang/CMakeLists.txt b/hipify-clang/CMakeLists.txt index 56de7c7904..09f6cb8a38 100644 --- a/hipify-clang/CMakeLists.txt +++ b/hipify-clang/CMakeLists.txt @@ -1,6 +1,11 @@ cmake_minimum_required(VERSION 2.8.12) project(hipify-clang) +if (MSVC AND MSVC_VERSION VERSION_LESS "1900") + message(SEND_ERROR "hipify-clang could be built by Visual Studio 14 2015 or higher.") + return() +endif() + find_package(LLVM REQUIRED) message(STATUS "Found LLVM ${LLVM_PACKAGE_VERSION}:") message(STATUS " - CMake module path: ${LLVM_CMAKE_DIR}") From 8a54ef365bf3c9226fbeef0bf4c56179546e3646 Mon Sep 17 00:00:00 2001 From: emankov Date: Fri, 4 Jan 2019 19:10:11 +0300 Subject: [PATCH 18/18] [HIPIFY][cmake] Change min supported version of cmake to 3.12.3 [Reason] CUDA 10 is supported by cmake since 3.12.3. --- hipify-clang/CMakeLists.txt | 2 +- hipify-clang/README.md | 6 +++--- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/hipify-clang/CMakeLists.txt b/hipify-clang/CMakeLists.txt index 09f6cb8a38..964e30d66a 100644 --- a/hipify-clang/CMakeLists.txt +++ b/hipify-clang/CMakeLists.txt @@ -1,4 +1,4 @@ -cmake_minimum_required(VERSION 2.8.12) +cmake_minimum_required(VERSION 3.12.3) project(hipify-clang) if (MSVC AND MSVC_VERSION VERSION_LESS "1900") diff --git a/hipify-clang/README.md b/hipify-clang/README.md index e553fb4adc..8ef9883b55 100644 --- a/hipify-clang/README.md +++ b/hipify-clang/README.md @@ -55,7 +55,7 @@ In most cases, you can get a suitable version of LLVM+CLANG with your package manager. Failing that or having multiple versions of LLVM, you can [download a release archive](http://releases.llvm.org/), build or install it, and set -[CMAKE_PREFIX_PATH](https://cmake.org/cmake/help/v3.10/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it; for instance: `-DCMAKE_PREFIX_PATH=f:\LLVM\6.0.1\dist` +[CMAKE_PREFIX_PATH](https://cmake.org/cmake/help/v3.12/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it; for instance: `-DCMAKE_PREFIX_PATH=f:\LLVM\6.0.1\dist` ## Build and install @@ -175,7 +175,7 @@ LLVM 5.0.0 - 6.0.1, CUDA 8.0, cudnn-8.0 Build system for the above configurations: -Python 2.7 (min), cmake 3.5.2 (min), GNU C/C++ 5.4.0 (min). +Python 2.7 (min), cmake 3.12.3 (min), GNU C/C++ 5.4.0 (min). Here is an example of building `hipify-clang` with testing support on `Ubuntu 16.04`: @@ -278,7 +278,7 @@ LLVM 5.0.0 - 5.0.2, CUDA 8.0, cudnn-8.0 Build system for the above configurations: -Python 3.6 (min), cmake 3.10 (min), Visual Studio 15.5 2017 (min). +Python 3.6 (min), cmake 3.12.3 (min), Visual Studio 15.5 2017 (min). Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 15 2017`: