From e9532108bb1674e7262616f8e3401e499c1f89d1 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 [ROCm/hip commit: ec03a8b3520235b085bfc915f1047846ad6a7f91] --- .../include/hip/hcc_detail/hip_runtime_api.h | 2 +- .../hip/tests/src/kernel/hipGridLaunch.cpp | 31 +++++++++++++++++++ 2 files changed, 32 insertions(+), 1 deletion(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index 97d6320f7f..d4b8c1faf4 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/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/projects/hip/tests/src/kernel/hipGridLaunch.cpp b/projects/hip/tests/src/kernel/hipGridLaunch.cpp index 8505a0758f..4f7883c656 100644 --- a/projects/hip/tests/src/kernel/hipGridLaunch.cpp +++ b/projects/hip/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 42a683e1594aa5948026f8af53bc9b0874d1e241 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 [ROCm/hip commit: 4272119eeb8e8d277e4b44f2c893de4dda8c350b] --- projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp b/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp index 53e21c93bb..151a86a6c0 100644 --- a/projects/hip/tests/src/runtimeApi/memory/hipMemset3D.cpp +++ b/projects/hip/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 279550644e7fb7aff2f3b6a59e7b2bea2bf1506a 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 [ROCm/hip commit: 208d7cfaf14c433cbc9a253db90db82f60591c89] --- projects/hip/hipify-clang/src/LLVMCompat.cpp | 9 +++++++++ projects/hip/hipify-clang/src/LLVMCompat.h | 3 +++ projects/hip/hipify-clang/src/main.cpp | 6 +++--- 3 files changed, 15 insertions(+), 3 deletions(-) diff --git a/projects/hip/hipify-clang/src/LLVMCompat.cpp b/projects/hip/hipify-clang/src/LLVMCompat.cpp index f96bd86bb0..22a0b56d64 100644 --- a/projects/hip/hipify-clang/src/LLVMCompat.cpp +++ b/projects/hip/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/projects/hip/hipify-clang/src/LLVMCompat.h b/projects/hip/hipify-clang/src/LLVMCompat.h index 4d6edcc9f8..98ff66fb41 100644 --- a/projects/hip/hipify-clang/src/LLVMCompat.h +++ b/projects/hip/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/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp index 06981a9d70..7dc42f8882 100644 --- a/projects/hip/hipify-clang/src/main.cpp +++ b/projects/hip/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 199a2c7189667de2848ae2025a3b2afd302b6c4f 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 [ROCm/hip commit: d5d6541ad5a9fd5dfa7b355c87eedd3164414d70] --- projects/hip/Jenkinsfile | 47 +--------------------------------------- 1 file changed, 1 insertion(+), 46 deletions(-) diff --git a/projects/hip/Jenkinsfile b/projects/hip/Jenkinsfile index 1257bfb0b4..812b68633a 100644 --- a/projects/hip/Jenkinsfile +++ b/projects/hip/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 d6593e4138b62313185b2e1cff88cb73e4f011f8 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. [ROCm/hip commit: 2d3f62c3e11668030a026f243982f3728c6c7b40] --- projects/hip/hipify-clang/src/main.cpp | 6 +++++- projects/hip/tests/hipify-clang/lit.cfg | 2 +- 2 files changed, 6 insertions(+), 2 deletions(-) diff --git a/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp index 7dc42f8882..8549e80d18 100644 --- a/projects/hip/hipify-clang/src/main.cpp +++ b/projects/hip/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/projects/hip/tests/hipify-clang/lit.cfg b/projects/hip/tests/hipify-clang/lit.cfg index d0328c06c2..e5805a5663 100644 --- a/projects/hip/tests/hipify-clang/lit.cfg +++ b/projects/hip/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 4e50eba1012226783ef8f42943487e8a493f5356 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. [ROCm/hip commit: 7ea586c323ee0c3df0f6d9d09937e215a31c023e] --- projects/hip/hipify-clang/README.md | 1 - projects/hip/hipify-clang/src/main.cpp | 2 ++ projects/hip/tests/hipify-clang/lit.cfg | 2 +- 3 files changed, 3 insertions(+), 2 deletions(-) diff --git a/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md index 2df3a73bff..845d0e0a9e 100644 --- a/projects/hip/hipify-clang/README.md +++ b/projects/hip/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/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp index 8549e80d18..a8f7e6a2f9 100644 --- a/projects/hip/hipify-clang/src/main.cpp +++ b/projects/hip/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/projects/hip/tests/hipify-clang/lit.cfg b/projects/hip/tests/hipify-clang/lit.cfg index e5805a5663..db1508a38f 100644 --- a/projects/hip/tests/hipify-clang/lit.cfg +++ b/projects/hip/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 142d00b382dd83c469bae345ecb02b0a72c0e99b 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. [ROCm/hip commit: dec459efcadb204161e00fdb9e185cc349364263] --- projects/hip/hipify-clang/README.md | 3 +-- projects/hip/tests/hipify-clang/lit.cfg | 2 +- 2 files changed, 2 insertions(+), 3 deletions(-) diff --git a/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md index 845d0e0a9e..5b4424b791 100644 --- a/projects/hip/hipify-clang/README.md +++ b/projects/hip/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/projects/hip/tests/hipify-clang/lit.cfg b/projects/hip/tests/hipify-clang/lit.cfg index db1508a38f..7b25731ba6 100644 --- a/projects/hip/tests/hipify-clang/lit.cfg +++ b/projects/hip/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 01a335b975e4e9478d97e7166ddc682208133045 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. [ROCm/hip commit: c519c89abffd5eaa50ee5ff4e29b004fc9782e6c] --- projects/hip/hipify-clang/src/main.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp index a8f7e6a2f9..d2d3d074b4 100644 --- a/projects/hip/hipify-clang/src/main.cpp +++ b/projects/hip/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 dc9f7851520bf00837cba6d9172a002bfa080ef1 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. [ROCm/hip commit: c09022232a7732bd70b4f39dd4adce8d539b3aa5] --- projects/hip/hipify-clang/src/main.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp index d2d3d074b4..2a3b86f204 100644 --- a/projects/hip/hipify-clang/src/main.cpp +++ b/projects/hip/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 a2ab2b0370af9bac5381271a9440a30a8a2afcf9 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 [ROCm/hip commit: 548f4dd4e51250a579d1b5b57154ea2368be1fca] --- projects/hip/include/hip/hcc_detail/hip_runtime_api.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index d4b8c1faf4..0d526ca25d 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/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 7bcb83a05f111b8657fb00d68b29ecd852f81250 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. [ROCm/hip commit: 25c7e5d6096330075559a1fea50a58ffb1a41295] --- .../hip/hcc_detail/functional_grid_launch.hpp | 38 +++++-- .../include/hip/hcc_detail/program_state.hpp | 2 + projects/hip/src/program_state.cpp | 104 ++++++++++++++++++ 3 files changed, 137 insertions(+), 7 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp b/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp index e678f25aa2..cd90aa401a 100644 --- a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/projects/hip/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/projects/hip/include/hip/hcc_detail/program_state.hpp b/projects/hip/include/hip/hcc_detail/program_state.hpp index bdb87b3509..92bef22172 100644 --- a/projects/hip/include/hip/hcc_detail/program_state.hpp +++ b/projects/hip/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/projects/hip/src/program_state.cpp b/projects/hip/src/program_state.cpp index 88cdeeb404..922d827be4 100644 --- a/projects/hip/src/program_state.cpp +++ b/projects/hip/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 e127990e23173d412e54e0d0291c102c9fca3b8e 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. [ROCm/hip commit: 340674ceb6ee471e50ed957b51d6a27847176c0f] --- projects/hip/src/program_state.cpp | 13 +++++++------ 1 file changed, 7 insertions(+), 6 deletions(-) diff --git a/projects/hip/src/program_state.cpp b/projects/hip/src/program_state.cpp index 922d827be4..38cae74dcc 100644 --- a/projects/hip/src/program_state.cpp +++ b/projects/hip/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 587745b841bd0f9726d75f271255b77f2562ef28 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. [ROCm/hip commit: ec14daa7cee0ea5c1073dfc3c369c5ab1a50b37d] --- .../hip/hcc_detail/functional_grid_launch.hpp | 21 +++++++++----- .../include/hip/hcc_detail/program_state.hpp | 5 ++-- projects/hip/src/program_state.cpp | 29 +++++++++++++++---- 3 files changed, 40 insertions(+), 15 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp b/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp index cd90aa401a..2fbda48629 100644 --- a/projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/projects/hip/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/projects/hip/include/hip/hcc_detail/program_state.hpp b/projects/hip/include/hip/hcc_detail/program_state.hpp index 92bef22172..da13c7c3db 100644 --- a/projects/hip/include/hip/hcc_detail/program_state.hpp +++ b/projects/hip/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/projects/hip/src/program_state.cpp b/projects/hip/src/program_state.cpp index 38cae74dcc..bb906b0ad9 100644 --- a/projects/hip/src/program_state.cpp +++ b/projects/hip/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 8883153c059393cae6f82d7275b4b3fbd49730cf 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 [ROCm/hip commit: 7fb0f4ee2a4a5a2ce933b99796054035d595bf43] --- .../tests/hipify-clang/unit_tests/samples/cudaRegister.cu | 8 -------- 1 file changed, 8 deletions(-) diff --git a/projects/hip/tests/hipify-clang/unit_tests/samples/cudaRegister.cu b/projects/hip/tests/hipify-clang/unit_tests/samples/cudaRegister.cu index 43b4345337..11192c452e 100644 --- a/projects/hip/tests/hipify-clang/unit_tests/samples/cudaRegister.cu +++ b/projects/hip/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) [ROCm/hip commit: 65155c7159b35c45fad5992ea73a2555a0c0b9cf] --- projects/hip/tests/hipify-clang/lit.cfg | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/projects/hip/tests/hipify-clang/lit.cfg b/projects/hip/tests/hipify-clang/lit.cfg index 7b25731ba6..462aba6c52 100644 --- a/projects/hip/tests/hipify-clang/lit.cfg +++ b/projects/hip/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 74f77c9f29940d0d8c37e47f9605ed5f392a6af2 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 [ROCm/hip commit: 9abcad9407d6c607a29516d349faf4d52ce7ed51] --- projects/hip/hipify-clang/README.md | 2 +- projects/hip/tests/hipify-clang/lit.cfg | 2 ++ projects/hip/tests/hipify-clang/lit.site.cfg.in | 1 + 3 files changed, 4 insertions(+), 1 deletion(-) diff --git a/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md index 5b4424b791..e553fb4adc 100644 --- a/projects/hip/hipify-clang/README.md +++ b/projects/hip/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/projects/hip/tests/hipify-clang/lit.cfg b/projects/hip/tests/hipify-clang/lit.cfg index 462aba6c52..76a811031a 100644 --- a/projects/hip/tests/hipify-clang/lit.cfg +++ b/projects/hip/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/projects/hip/tests/hipify-clang/lit.site.cfg.in b/projects/hip/tests/hipify-clang/lit.site.cfg.in index f7302252a8..f93c206be4 100644 --- a/projects/hip/tests/hipify-clang/lit.site.cfg.in +++ b/projects/hip/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 e371a454008c4969492671b12d280f3c94e31fae 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". [ROCm/hip commit: cd15a856638d0a2dcc1c5c999d3baca37cdcafba] --- projects/hip/hipify-clang/CMakeLists.txt | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/projects/hip/hipify-clang/CMakeLists.txt b/projects/hip/hipify-clang/CMakeLists.txt index 56de7c7904..09f6cb8a38 100644 --- a/projects/hip/hipify-clang/CMakeLists.txt +++ b/projects/hip/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 3fa069c87d5ce0dc2f3ea4f6f7672a42c1f9c60f 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. [ROCm/hip commit: 8a54ef365bf3c9226fbeef0bf4c56179546e3646] --- projects/hip/hipify-clang/CMakeLists.txt | 2 +- projects/hip/hipify-clang/README.md | 6 +++--- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/projects/hip/hipify-clang/CMakeLists.txt b/projects/hip/hipify-clang/CMakeLists.txt index 09f6cb8a38..964e30d66a 100644 --- a/projects/hip/hipify-clang/CMakeLists.txt +++ b/projects/hip/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/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md index e553fb4adc..8ef9883b55 100644 --- a/projects/hip/hipify-clang/README.md +++ b/projects/hip/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`: