diff --git a/hip-config-clang.cmake.in b/hip-config-clang.cmake.in index 199dca774a..9680b2ff8e 100644 --- a/hip-config-clang.cmake.in +++ b/hip-config-clang.cmake.in @@ -1,16 +1,4 @@ @PACKAGE_INIT@ -if(NOT DEFINED HSA_PATH) - if(NOT DEFINED ENV{HSA_PATH}) - set(HSA_PATH "/opt/rocm/hsa" CACHE PATH "Path to which HSA runtime has been installed") - else() - set(HSA_PATH $ENV{HSA_PATH} CACHE PATH "Path to which HSA runtime has been installed") - endif() -endif() -if(IS_ABSOLUTE ${HSA_PATH} AND EXISTS ${HSA_PATH} AND IS_DIRECTORY ${HSA_PATH}) - message(STATUS "Looking for HSA runtime in: " ${HSA_PATH}) -else() - message(FATAL_ERROR "Don't know where to find HSA runtime. Please specify absolute path using -DHSA_PATH") -endif() include(CMakeFindDependencyMacro OPTIONAL RESULT_VARIABLE _CMakeFindDependencyMacro_FOUND) if (NOT _CMakeFindDependencyMacro_FOUND) @@ -55,8 +43,7 @@ set(HIP_COMPILER "@HIP_COMPILER@") set(HIP_RUNTIME "@HIP_RUNTIME@") set_and_check( hip_INCLUDE_DIR "@PACKAGE_INCLUDE_INSTALL_DIR@" ) -set_and_check( HSA_INCLUDE_PATH "${HSA_PATH}/include" ) -set( hip_INCLUDE_DIRS "${hip_INCLUDE_DIR}" "${HSA_INCLUDE_PATH}" ) +set_and_check( hip_INCLUDE_DIRS "${hip_INCLUDE_DIR}" ) set_and_check( hip_LIB_INSTALL_DIR "@PACKAGE_LIB_INSTALL_DIR@" ) set_and_check( hip_BIN_INSTALL_DIR "@PACKAGE_BIN_INSTALL_DIR@" ) diff --git a/hip-config-hcc.cmake.in b/hip-config-hcc.cmake.in index b385f5ea8e..c0ffc6e2af 100644 --- a/hip-config-hcc.cmake.in +++ b/hip-config-hcc.cmake.in @@ -1,16 +1,4 @@ @PACKAGE_INIT@ -if(NOT DEFINED HSA_PATH) - if(NOT DEFINED ENV{HSA_PATH}) - set(HSA_PATH "/opt/rocm/hsa" CACHE PATH "Path to which HSA runtime has been installed") - else() - set(HSA_PATH $ENV{HSA_PATH} CACHE PATH "Path to which HSA runtime has been installed") - endif() -endif() -if(IS_ABSOLUTE ${HSA_PATH} AND EXISTS ${HSA_PATH} AND IS_DIRECTORY ${HSA_PATH}) - message(STATUS "Looking for HSA runtime in: " ${HSA_PATH}) -else() - message(FATAL_ERROR "Don't know where to find HSA runtime. Please specify absolute path using -DHSA_PATH") -endif() include(CMakeFindDependencyMacro OPTIONAL RESULT_VARIABLE _CMakeFindDependencyMacro_FOUND) if (NOT _CMakeFindDependencyMacro_FOUND) @@ -55,8 +43,7 @@ set(HIP_COMPILER "@HIP_COMPILER@") set(HIP_RUNTIME "@HIP_RUNTIME@") set_and_check( hip_INCLUDE_DIR "@PACKAGE_INCLUDE_INSTALL_DIR@" ) -set_and_check( HSA_INCLUDE_PATH "${HSA_PATH}/include" ) -set( hip_INCLUDE_DIRS "${hip_INCLUDE_DIR}" "${HSA_INCLUDE_PATH}" ) +set_and_check( hip_INCLUDE_DIRS "${hip_INCLUDE_DIR}" ) set_and_check( hip_LIB_INSTALL_DIR "@PACKAGE_LIB_INSTALL_DIR@" ) set_and_check( hip_BIN_INSTALL_DIR "@PACKAGE_BIN_INSTALL_DIR@" ) diff --git a/hip_prof_gen.py b/hip_prof_gen.py index 587b2d545a..e2876d87bc 100755 --- a/hip_prof_gen.py +++ b/hip_prof_gen.py @@ -333,7 +333,7 @@ def generate_prof_header(f, api_map, opts_map): # Generating the callbacks ID enumaration f.write('\n// Return HIP API string\n') - f.write('static inline const char* hip_api_name(const uint32_t& id) {\n') + f.write('static inline const char* hip_api_name(const uint32_t id) {\n') f.write(' switch(id) {\n') for name in api_map.keys(): f.write(' case HIP_API_ID_' + name + ': return "' + name + '";\n') @@ -353,7 +353,10 @@ def generate_prof_header(f, api_map, opts_map): if len(args) != 0: f.write(' struct {\n') for arg_tuple in args: - f.write(' ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n') + if arg_tuple[0] == "hipLimit_t": + f.write(' enum ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n') + else: + f.write(' ' + arg_tuple[0] + ' ' + arg_tuple[1] + ';\n') f.write(' } ' + name + ';\n') f.write( ' } args;\n' + diff --git a/hipify-clang/README.md b/hipify-clang/README.md index 470e5db257..670e8206ee 100644 --- a/hipify-clang/README.md +++ b/hipify-clang/README.md @@ -143,9 +143,9 @@ To run it: * Path to cuDNN should be specified by the `CUDA_DNN_ROOT_DIR` option: - - Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.0-v7.6.1.34` + - Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.0-v7.6.2.24` - - Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.6.1.34` + - Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.6.2.24` 5. Ensure [`python`](https://www.python.org/downloads) of minimum required version 2.7 is installed. @@ -179,9 +179,9 @@ To run it: On Linux the following configurations are tested: -Ubuntu 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-5.0.5 - cudnn-7.6.1.34 +Ubuntu 14: LLVM 5.0.0 - 6.0.1, CUDA 7.0 - 9.0, cudnn-5.0.5 - cudnn-7.6.2.24 -Ubuntu 16-18: LLVM 8.0.0 - 8.0.1, CUDA 8.0 - 10.0, cudnn-5.1.10 - cudnn-7.6.1.34 +Ubuntu 16-18: LLVM 8.0.0 - 8.0.1, CUDA 8.0 - 10.0, cudnn-5.1.10 - cudnn-7.6.2.24 Build system for the above configurations: @@ -196,7 +196,7 @@ cmake -DCMAKE_INSTALL_PREFIX=../dist \ -DCMAKE_PREFIX_PATH=/srv/git/LLVM/8.0.1/dist \ -DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-10.0 \ - -DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.0-v7.6.1.34 \ + -DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.0-v7.6.2.24 \ -DLLVM_EXTERNAL_LIT=/srv/git/LLVM/8.0.1/build/bin/llvm-lit \ .. ``` @@ -248,65 +248,67 @@ Running HIPify regression tests CUDA 10.0 - will be used for testing LLVM 8.0.1 - will be used for testing x86_64 - Platform architecture -Linux 4.13.0-32-generic - Platform OS +Linux 5.2.0-rc1-kfd-compute-roc-master-int-hipclang-623 - Platform OS 64 - hipify-clang binary bitness 64 - python 2.7.12 binary bitness ======================================== --- Testing: 52 tests, 12 threads -- -PASS: hipify :: unit_tests/headers/headers_test_03.cu (1 of 52) -PASS: hipify :: unit_tests/headers/headers_test_02.cu (2 of 52) -PASS: hipify :: unit_tests/headers/headers_test_01.cu (3 of 52) -PASS: hipify :: unit_tests/headers/headers_test_05.cu (4 of 52) -PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (5 of 52) -PASS: hipify :: unit_tests/headers/headers_test_07.cu (6 of 52) -PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (7 of 52) -PASS: hipify :: unit_tests/headers/headers_test_06.cu (8 of 52) -PASS: hipify :: unit_tests/headers/headers_test_04.cu (9 of 52) -PASS: hipify :: unit_tests/headers/headers_test_08.cu (10 of 52) -PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (11 of 52) -PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (12 of 52) -PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (13 of 52) -PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (14 of 52) -PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (15 of 52) -PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (16 of 52) -PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (17 of 52) -PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (18 of 52) -PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (19 of 52) -PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (20 of 52) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (21 of 52) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (22 of 52) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (23 of 52) -PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (24 of 52) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (25 of 52) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (26 of 52) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (27 of 52) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (28 of 52) -PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp (29 of 52) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (30 of 52) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (31 of 52) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (32 of 52) -PASS: hipify :: unit_tests/headers/headers_test_09.cu (33 of 52) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (34 of 52) -PASS: hipify :: unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp (35 of 52) -PASS: hipify :: unit_tests/samples/allocators.cu (36 of 52) -PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp (37 of 52) -PASS: hipify :: unit_tests/samples/coalescing.cu (38 of 52) -PASS: hipify :: unit_tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp (39 of 52) -PASS: hipify :: unit_tests/samples/2_Cookbook/2_Profiler/Profiler.cpp (40 of 52) -PASS: hipify :: unit_tests/samples/2_Cookbook/7_streams/stream.cpp (41 of 52) -PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (42 of 52) -PASS: hipify :: unit_tests/samples/axpy.cu (43 of 52) -PASS: hipify :: unit_tests/samples/2_Cookbook/8_peer2peer/peer2peer.cpp (44 of 52) -PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (45 of 52) -PASS: hipify :: unit_tests/samples/intro.cu (46 of 52) -PASS: hipify :: unit_tests/samples/square.cu (47 of 52) -PASS: hipify :: unit_tests/samples/cudaRegister.cu (48 of 52) -PASS: hipify :: unit_tests/samples/vec_add.cu (49 of 52) -PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu (50 of 52) -PASS: hipify :: unit_tests/samples/static_shared_memory.cu (51 of 52) -PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (52 of 52) -Testing Time: 3.05s - Expected Passes : 52 +-- Testing: 54 tests, 12 threads -- +PASS: hipify :: unit_tests/headers/headers_test_03.cu (1 of 54) +PASS: hipify :: unit_tests/headers/headers_test_02.cu (2 of 54) +PASS: hipify :: unit_tests/headers/headers_test_01.cu (3 of 54) +PASS: hipify :: unit_tests/headers/headers_test_05.cu (4 of 54) +PASS: hipify :: unit_tests/headers/headers_test_11.cu (5 of 54) +PASS: hipify :: unit_tests/headers/headers_test_10.cu (6 of 54) +PASS: hipify :: unit_tests/headers/headers_test_07.cu (7 of 54) +PASS: hipify :: unit_tests/headers/headers_test_06.cu (8 of 54) +PASS: hipify :: unit_tests/headers/headers_test_04.cu (9 of 54) +PASS: hipify :: unit_tests/headers/headers_test_08.cu (10 of 54) +PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (11 of 54) +PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (12 of 54) +PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (13 of 54) +PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (14 of 54) +PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (15 of 54) +PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (16 of 54) +PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (17 of 54) +PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (18 of 54) +PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (19 of 54) +PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (20 of 54) +PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (21 of 54) +PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (22 of 54) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (23 of 54) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (24 of 54) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (25 of 54) +PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (26 of 54) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (27 of 54) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (28 of 54) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (29 of 54) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (30 of 54) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (31 of 54) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (32 of 54) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (33 of 54) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (34 of 54) +PASS: hipify :: unit_tests/headers/headers_test_09.cu (35 of 54) +PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp (36 of 54) +PASS: hipify :: unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp (37 of 54) +PASS: hipify :: unit_tests/samples/allocators.cu (38 of 54) +PASS: hipify :: unit_tests/samples/coalescing.cu (39 of 54) +PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp (40 of 54) +PASS: hipify :: unit_tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp (41 of 54) +PASS: hipify :: unit_tests/samples/2_Cookbook/2_Profiler/Profiler.cpp (42 of 54) +PASS: hipify :: unit_tests/samples/2_Cookbook/7_streams/stream.cpp (43 of 54) +PASS: hipify :: unit_tests/samples/2_Cookbook/8_peer2peer/peer2peer.cpp (44 of 54) +PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (45 of 54) +PASS: hipify :: unit_tests/samples/axpy.cu (46 of 54) +PASS: hipify :: unit_tests/samples/intro.cu (47 of 54) +PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (48 of 54) +PASS: hipify :: unit_tests/samples/cudaRegister.cu (49 of 54) +PASS: hipify :: unit_tests/samples/vec_add.cu (50 of 54) +PASS: hipify :: unit_tests/samples/square.cu (51 of 54) +PASS: hipify :: unit_tests/samples/static_shared_memory.cu (52 of 54) +PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu (53 of 54) +PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (54 of 54) +Testing Time: 2.92s + Expected Passes : 54 [100%] Built target test-hipify ``` ### Windows @@ -315,13 +317,13 @@ On Windows 10 the following configurations are tested: LLVM 5.0.0 - 5.0.2, CUDA 8.0, cudnn-5.1.10 - cudnn-7.1.4.18 -LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn-7.0.5.15 - cudnn-7.6.1.34 +LLVM 6.0.0 - 6.0.1, CUDA 9.0, cudnn-7.0.5.15 - cudnn-7.6.2.24 -LLVM 7.0.0 - 8.0.1 (with patch*), CUDA 7.5 - 10.0, cudnn-7.0.5.15 - cudnn-7.6.1.34 +LLVM 7.0.0 - 8.0.1 (with patch*), CUDA 7.5 - 10.0, cudnn-7.0.5.15 - cudnn-7.6.2.24 Build system for the above configurations: -Python 3.6 (min), cmake 3.12.3 (min), Visual Studio 15.5.2 2017 - 16.1.6 2019. +Python 3.6 (min), cmake 3.12.3 (min), Visual Studio 2017 (15.5.2) - 2019 (16.2.0). Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 15 2017`: @@ -334,7 +336,7 @@ cmake -DCMAKE_PREFIX_PATH=f:/LLVM/6.0.1/dist \ -DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0" \ -DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v9.0" \ - -DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.6.1.34 \ + -DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.6.2.24 \ -DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.1/build/Release/bin/llvm-lit.py \ -Thost=x64 .. @@ -345,7 +347,7 @@ cmake -- - CMake module path: F:/LLVM/6.0.1/dist/lib/cmake/llvm -- - Include path : F:/LLVM/6.0.1/dist/include -- - Binary path : F:/LLVM/6.0.1/dist/bin --- Found PythonInterp: C:/Program Files/Python36/python.exe (found suitable version "3.6.4", minimum required is "3.6") +-- Found PythonInterp: C:/Program Files/Python36/python.exe (found suitable version "3.7.4", minimum required is "3.6") -- Found lit: C:/Program Files/Python36/Scripts/lit.exe -- Found FileCheck: F:/LLVM/6.0.1/dist/bin/FileCheck.exe -- Found CUDA: C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0 (found version "9.0") diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp index d6fe9b88e5..52ba2eeaaf 100644 --- a/hipify-clang/src/HipifyAction.cpp +++ b/hipify-clang/src/HipifyAction.cpp @@ -25,6 +25,7 @@ THE SOFTWARE. #include "clang/Frontend/CompilerInstance.h" #include "clang/ASTMatchers/ASTMatchFinder.h" #include "clang/ASTMatchers/ASTMatchers.h" +#include "clang/Lex/HeaderSearch.h" #include "LLVMCompat.h" #include "CUDA2HIP.h" #include "StringUtils.h" @@ -398,20 +399,50 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi return Finder->newASTConsumer(); } +void HipifyAction::Ifndef(clang::SourceLocation Loc, const clang::Token &MacroNameTok, const clang::MacroDefinition &MD) { + clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + if (!SM.isWrittenInMainFile(Loc)) { + return; + } + StringRef Text(SM.getCharacterData(MacroNameTok.getLocation()), MacroNameTok.getLength()); + Ifndefs.insert(std::make_pair(Text.str(), MacroNameTok.getEndLoc())); +} + void HipifyAction::EndSourceFileAction() { // Insert the hip header, if we didn't already do it by accident during substitution. if (!insertedRuntimeHeader) { // It's not sufficient to just replace CUDA headers with hip ones, because numerous CUDA headers are // implicitly included by the compiler. Instead, we _delete_ CUDA headers, and unconditionally insert // one copy of the hip include into every file. + bool placeForIncludeCalculated = false; + clang::SourceLocation sl, controllingMacroLoc; clang::SourceManager& SM = getCompilerInstance().getSourceManager(); - clang::SourceLocation sl; + clang::Preprocessor& PP = getCompilerInstance().getPreprocessor(); + clang::HeaderSearch& HS = PP.getHeaderSearchInfo(); + clang::ExternalPreprocessorSource* EPL = HS.getExternalLookup(); + const clang::FileEntry* FE = SM.getFileEntryForID(SM.getMainFileID()); + const clang::IdentifierInfo* controllingMacro = HS.getFileInfo(FE).getControllingMacro(EPL); + if (controllingMacro) { + auto found = Ifndefs.find(controllingMacro->getName().str()); + if (found != Ifndefs.end()) { + controllingMacroLoc = found->second; + placeForIncludeCalculated = true; + } + } if (pragmaOnce) { - sl = pragmaOnceLoc; - } else if (firstHeader) { - sl = firstHeaderLoc; - } else { - sl = SM.getLocForStartOfFile(SM.getMainFileID()); + if (placeForIncludeCalculated) { + sl = pragmaOnceLoc < controllingMacroLoc ? pragmaOnceLoc : controllingMacroLoc; + } else { + sl = pragmaOnceLoc; + } + placeForIncludeCalculated = true; + } + if (!placeForIncludeCalculated) { + if (firstHeader) { + sl = firstHeaderLoc; + } else { + sl = SM.getLocForStartOfFile(SM.getMainFileID()); + } } clang::FullSourceLoc fullSL(sl, SM); ct::Replacement Rep(SM, sl, 0, "\n#include \n"); @@ -445,6 +476,10 @@ public: void PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer) override { hipifyAction.PragmaDirective(Loc, Introducer); } + + void Ifndef(clang::SourceLocation Loc, const clang::Token &MacroNameTok, const clang::MacroDefinition &MD) override { + hipifyAction.Ifndef(Loc, MacroNameTok, MD); + } }; } diff --git a/hipify-clang/src/HipifyAction.h b/hipify-clang/src/HipifyAction.h index cc7f3799ba..adafba2df9 100644 --- a/hipify-clang/src/HipifyAction.h +++ b/hipify-clang/src/HipifyAction.h @@ -40,6 +40,7 @@ class HipifyAction : public clang::ASTFrontendAction, public clang::ast_matchers::MatchFinder::MatchCallback { private: ct::Replacements* replacements; + std::map Ifndefs; std::unique_ptr Finder; // CUDA implicitly adds its runtime header. We rewrite explicitly-provided CUDA includes with equivalent // ones, and track - using this flag - if the result led to us including the hip runtime header. If it did @@ -81,6 +82,9 @@ public: const clang::Module *imported); // Called by the preprocessor for each pragma directive during the non-raw lexing pass. void PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer); + // Called by the preprocessor for each ifndef directive during the non-raw lexing pass. + // Found ifndef will be used in EndSourceFileAction() for catching include guard controlling macro. + void Ifndef(clang::SourceLocation Loc, const clang::Token &MacroNameTok, const clang::MacroDefinition &MD); protected: // Add a Replacement for the current file. These will all be applied after executing the FrontendAction. diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index d275bb1820..465655680f 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -736,21 +736,13 @@ int __any(int predicate) { __device__ inline unsigned long long int __ballot(int predicate) { -#if defined(__HCC__) - return __llvm_amdgcn_icmp_i32(predicate, 0, ICMP_NE); -#else - return __builtin_amdgcn_uicmp(predicate, 0, ICMP_NE); -#endif + return __builtin_amdgcn_uicmp(predicate, 0, ICMP_NE); } __device__ inline unsigned long long int __ballot64(int predicate) { -#if defined(__HCC__) - return __llvm_amdgcn_icmp_i32(predicate, 0, ICMP_NE); -#else - return __builtin_amdgcn_uicmp(predicate, 0, ICMP_NE); -#endif + return __builtin_amdgcn_uicmp(predicate, 0, ICMP_NE); } // hip.amdgcn.bc - lanemask diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 7e887d1172..e8871912fa 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -45,7 +45,7 @@ THE SOFTWARE. #include #include -#if !__HIP_VDI__ +#if !__HIP_VDI__ && defined(__cplusplus) #include #include #endif @@ -79,9 +79,11 @@ THE SOFTWARE. #define __dparm(x) #endif +#ifdef __cplusplus namespace hip_impl { hipError_t hip_init(); } // namespace hip_impl +#endif // Structure definitions: #ifdef __cplusplus @@ -1459,6 +1461,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, #else hipError_t hipModuleGetGlobal(void**, size_t*, hipModule_t, const char*); +#ifdef __cplusplus //Start : Not supported in gcc namespace hip_impl { inline __attribute__((visibility("hidden"))) @@ -1466,6 +1469,7 @@ hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes, const char* name); } // Namespace hip_impl. + /** * @brief Copies the memory address of symbol @p symbolName to @p devPtr * @@ -1504,15 +1508,18 @@ hipError_t hipGetSymbolSize(size_t* size, const void* symbolName) { void* devPtr = nullptr; return hip_impl::read_agent_global_from_process(&devPtr, size, (const char*)symbolName); } +#endif // End : Not supported in gcc #if defined(__cplusplus) } // extern "C" #endif +#ifdef __cplusplus namespace hip_impl { hipError_t hipMemcpyToSymbol(void*, const void*, size_t, size_t, hipMemcpyKind, const char*); } // Namespace hip_impl. +#endif #if defined(__cplusplus) extern "C" { @@ -1541,6 +1548,7 @@ extern "C" { * hipMemcpyFromArrayAsync, hipMemcpy2DFromArrayAsync, hipMemcpyToSymbolAsync, * hipMemcpyFromSymbolAsync */ +#ifdef __cplusplus inline __attribute__((visibility("hidden"))) hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, @@ -1554,11 +1562,13 @@ hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, return hip_impl::hipMemcpyToSymbol(dst, src, sizeBytes, offset, kind, (const char*)symbolName); } +#endif #if defined(__cplusplus) } // extern "C" #endif +#ifdef __cplusplus namespace hip_impl { hipError_t hipMemcpyToSymbolAsync(void*, const void*, size_t, size_t, hipMemcpyKind, hipStream_t, const char*); @@ -1567,6 +1577,7 @@ hipError_t hipMemcpyFromSymbol(void*, const void*, size_t, size_t, hipError_t hipMemcpyFromSymbolAsync(void*, const void*, size_t, size_t, hipMemcpyKind, hipStream_t, const char*); } // Namespace hip_impl. +#endif #if defined(__cplusplus) extern "C" { @@ -1597,6 +1608,8 @@ extern "C" { * hipMemcpyFromArrayAsync, hipMemcpy2DFromArrayAsync, hipMemcpyToSymbolAsync, * hipMemcpyFromSymbolAsync */ + +#ifdef __cplusplus //Start : Not supported in gcc inline __attribute__((visibility("hidden"))) hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, @@ -1641,6 +1654,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, stream, (const char*)symbolName); } +#endif // End : Not supported in gcc #endif // __HIP_VDI__ /** @@ -2604,6 +2618,7 @@ hipError_t hipFuncGetAttributes(struct hipFuncAttributes* attr, const void* func } // extern "C" #endif +#ifdef __cplusplus namespace hip_impl { class agent_globals_impl; class agent_globals { @@ -2635,6 +2650,7 @@ namespace hip_impl { return get_agent_globals().read_agent_global_from_process(dptr, bytes, name); } } // Namespace hip_impl. +#endif #if defined(__cplusplus) extern "C" { @@ -3007,10 +3023,12 @@ const char* hipKernelNameRef(const hipFunction_t f); #ifdef __cplusplus +class TlsData; + hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* devPtr, const hipChannelFormatDesc* desc, size_t size = UINT_MAX); -hipError_t ihipBindTextureImpl(int dim, enum hipTextureReadMode readMode, size_t* offset, +hipError_t ihipBindTextureImpl(TlsData *tls, int dim, enum hipTextureReadMode readMode, size_t* offset, const void* devPtr, const struct hipChannelFormatDesc* desc, size_t size, textureReference* tex); @@ -3032,7 +3050,7 @@ hipError_t ihipBindTextureImpl(int dim, enum hipTextureReadMode readMode, size_t template hipError_t hipBindTexture(size_t* offset, struct texture& tex, const void* devPtr, const struct hipChannelFormatDesc& desc, size_t size = UINT_MAX) { - return ihipBindTextureImpl(dim, readMode, offset, devPtr, &desc, size, &tex); + return ihipBindTextureImpl(nullptr, dim, readMode, offset, devPtr, &desc, size, &tex); } /* @@ -3052,7 +3070,7 @@ hipError_t hipBindTexture(size_t* offset, struct texture& tex, template hipError_t hipBindTexture(size_t* offset, struct texture& tex, const void* devPtr, size_t size = UINT_MAX) { - return ihipBindTextureImpl(dim, readMode, offset, devPtr, &(tex.channelDesc), size, &tex); + return ihipBindTextureImpl(nullptr, dim, readMode, offset, devPtr, &(tex.channelDesc), size, &tex); } // C API @@ -3082,27 +3100,27 @@ hipError_t hipBindTexture2D(size_t* offset, struct texture& te hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array, const hipChannelFormatDesc* desc); -hipError_t ihipBindTextureToArrayImpl(int dim, enum hipTextureReadMode readMode, +hipError_t ihipBindTextureToArrayImpl(TlsData *tls, int dim, enum hipTextureReadMode readMode, hipArray_const_t array, const struct hipChannelFormatDesc& desc, textureReference* tex); template hipError_t hipBindTextureToArray(struct texture& tex, hipArray_const_t array) { - return ihipBindTextureToArrayImpl(dim, readMode, array, tex.channelDesc, &tex); + return ihipBindTextureToArrayImpl(nullptr, dim, readMode, array, tex.channelDesc, &tex); } template hipError_t hipBindTextureToArray(struct texture& tex, hipArray_const_t array, const struct hipChannelFormatDesc& desc) { - return ihipBindTextureToArrayImpl(dim, readMode, array, desc, &tex); + return ihipBindTextureToArrayImpl(nullptr, dim, readMode, array, desc, &tex); } template inline static hipError_t hipBindTextureToArray(struct texture *tex, hipArray_const_t array, const struct hipChannelFormatDesc* desc) { - return ihipBindTextureToArrayImpl(dim, readMode, array, *desc, tex); + return ihipBindTextureToArrayImpl(nullptr, dim, readMode, array, *desc, tex); } // C API diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h index 3407b8a752..a03a46b8cf 100644 --- a/include/hip/hcc_detail/hip_vector_types.h +++ b/include/hip/hcc_detail/hip_vector_types.h @@ -849,6 +849,7 @@ __MAKE_VECTOR_TYPE__(longlong, long long); __MAKE_VECTOR_TYPE__(float, float); __MAKE_VECTOR_TYPE__(double, double); +#ifdef __cplusplus #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \ static inline __device__ __host__ \ type make_##type(comp x) { type r{x}; return r; } @@ -867,6 +868,26 @@ __MAKE_VECTOR_TYPE__(double, double); type r{x, y, z, w}; \ return r; \ } +#else + #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \ + static inline __device__ __host__ \ + type make_##type(comp x) { type r; r.x =x; return r; } + + #define DECLOP_MAKE_TWO_COMPONENT(comp, type) \ + static inline __device__ __host__ \ + type make_##type(comp x, comp y) { type r; r.x=x; r.y=y; return r; } + + #define DECLOP_MAKE_THREE_COMPONENT(comp, type) \ + static inline __device__ __host__ \ + type make_##type(comp x, comp y, comp z) { type r; r.x=x; r.y=y; r.z=z; return r; } + + #define DECLOP_MAKE_FOUR_COMPONENT(comp, type) \ + static inline __device__ __host__ \ + type make_##type(comp x, comp y, comp z, comp w) { \ + type r; r.x=x; r.y=y; r.z=z; r.w=w; \ + return r; \ + } +#endif DECLOP_MAKE_ONE_COMPONENT(unsigned char, uchar1); DECLOP_MAKE_TWO_COMPONENT(unsigned char, uchar2); diff --git a/include/hip/hip_runtime_api.h b/include/hip/hip_runtime_api.h index c067f5bf54..1c8e971545 100644 --- a/include/hip/hip_runtime_api.h +++ b/include/hip/hip_runtime_api.h @@ -255,6 +255,8 @@ typedef enum __HIP_NODISCARD hipError_t { 1071, ///< Produced when the IPC memory attach failed from ROCr. hipErrorAssert = 1081, ///< Produced when the kernel calls assert. + hipErrorNotSupported = + 1082, ///< Produced when the hip API is not supported/implemented hipErrorTbd ///< Marker that more error codes are needed. } hipError_t; diff --git a/src/functional_grid_launch.inl b/src/functional_grid_launch.inl index 13679dcab3..3c611ed468 100644 --- a/src/functional_grid_launch.inl +++ b/src/functional_grid_launch.inl @@ -47,8 +47,8 @@ namespace hip_impl return *static_cast( stream->locked_getAv()->get_hsa_agent()); } - else if ( - ihipGetTlsDefaultCtx() && ihipGetTlsDefaultCtx()->getDevice()) { + GET_TLS(); + if (ihipGetTlsDefaultCtx() && ihipGetTlsDefaultCtx()->getDevice()) { return ihipGetDevice( ihipGetTlsDefaultCtx()->getDevice()->_deviceId)->_hsaAgent; } diff --git a/src/hip_clang.cpp b/src/hip_clang.cpp index b8beb05400..85dfb09f62 100644 --- a/src/hip_clang.cpp +++ b/src/hip_clang.cpp @@ -177,6 +177,7 @@ hipError_t hipConfigureCall( size_t sharedMem, hipStream_t stream) { + GET_TLS(); auto ctx = ihipGetTlsDefaultCtx(); LockedAccessor_CtxCrit_t crit(ctx->criticalData()); diff --git a/src/hip_context.cpp b/src/hip_context.cpp index e85c406cfc..300877f8b5 100644 --- a/src/hip_context.cpp +++ b/src/hip_context.cpp @@ -29,13 +29,10 @@ THE SOFTWARE. #include "hip_hcc_internal.h" #include "trace_helper.h" -// Stack of contexts -thread_local std::stack tls_ctxStack; -thread_local bool tls_getPrimaryCtx = true; - void ihipCtxStackUpdate() { - if (tls_ctxStack.empty()) { - tls_ctxStack.push(ihipGetTlsDefaultCtx()); + GET_TLS(); + if (tls->ctxStack.empty()) { + tls->ctxStack.push(ihipGetTlsDefaultCtx()); } } @@ -62,8 +59,8 @@ hipError_t hipCtxCreate(hipCtx_t* ctx, unsigned int flags, hipDevice_t device) { auto ictx = new ihipCtx_t(deviceHandle, g_deviceCnt, flags); *ctx = ictx; ihipSetTlsDefaultCtx(*ctx); - tls_ctxStack.push(*ctx); - tls_getPrimaryCtx = false; + tls->ctxStack.push(*ctx); + tls->getPrimaryCtx = false; deviceCrit->addContext(ictx); } @@ -119,7 +116,7 @@ hipError_t hipCtxDestroy(hipCtx_t ctx) { } else { if (currentCtx == ctx) { // need to destroy the ctx associated with calling thread - tls_ctxStack.pop(); + tls->ctxStack.pop(); } { auto deviceHandle = ctx->getWriteableDevice(); @@ -140,12 +137,12 @@ hipError_t hipCtxPopCurrent(hipCtx_t* ctx) { auto deviceHandle = currentCtx->getDevice(); *ctx = currentCtx; - if (!tls_ctxStack.empty()) { - tls_ctxStack.pop(); + if (!tls->ctxStack.empty()) { + tls->ctxStack.pop(); } - if (!tls_ctxStack.empty()) { - currentCtx = tls_ctxStack.top(); + if (!tls->ctxStack.empty()) { + currentCtx = tls->ctxStack.top(); } else { currentCtx = deviceHandle->_primaryCtx; } @@ -159,8 +156,8 @@ hipError_t hipCtxPushCurrent(hipCtx_t ctx) { hipError_t e = hipSuccess; if (ctx != NULL) { // TODO- is this check needed? ihipSetTlsDefaultCtx(ctx); - tls_ctxStack.push(ctx); - tls_getPrimaryCtx = false; + tls->ctxStack.push(ctx); + tls->getPrimaryCtx = false; } else { e = hipErrorInvalidContext; } @@ -170,10 +167,10 @@ hipError_t hipCtxPushCurrent(hipCtx_t ctx) { hipError_t hipCtxGetCurrent(hipCtx_t* ctx) { HIP_INIT_API(hipCtxGetCurrent, ctx); hipError_t e = hipSuccess; - if ((tls_getPrimaryCtx) || tls_ctxStack.empty()) { + if ((tls->getPrimaryCtx) || tls->ctxStack.empty()) { *ctx = ihipGetTlsDefaultCtx(); } else { - *ctx = tls_ctxStack.top(); + *ctx = tls->ctxStack.top(); } return ihipLogStatus(e); } @@ -182,11 +179,11 @@ hipError_t hipCtxSetCurrent(hipCtx_t ctx) { HIP_INIT_API(hipCtxSetCurrent, ctx); hipError_t e = hipSuccess; if (ctx == NULL) { - tls_ctxStack.pop(); + tls->ctxStack.pop(); } else { ihipSetTlsDefaultCtx(ctx); - tls_ctxStack.push(ctx); - tls_getPrimaryCtx = false; + tls->ctxStack.push(ctx); + tls->getPrimaryCtx = false; } return ihipLogStatus(e); } @@ -251,7 +248,7 @@ hipError_t hipCtxGetSharedMemConfig(hipSharedMemConfig* pConfig) { hipError_t hipCtxSynchronize(void) { HIP_INIT_API(hipCtxSynchronize, 1); - return ihipLogStatus(ihipSynchronize()); // TODP Shall check validity of ctx? + return ihipLogStatus(ihipSynchronize(tls)); // TODO Shall check validity of ctx? } hipError_t hipCtxGetFlags(unsigned int* flags) { diff --git a/src/hip_device.cpp b/src/hip_device.cpp index 1c0e5c9109..2d74985b64 100644 --- a/src/hip_device.cpp +++ b/src/hip_device.cpp @@ -136,14 +136,14 @@ hipError_t hipSetDevice(int deviceId) { return ihipLogStatus(hipErrorInvalidDevice); } else { ihipSetTlsDefaultCtx(ihipGetPrimaryCtx(deviceId)); - tls_getPrimaryCtx = true; + tls->getPrimaryCtx = true; return ihipLogStatus(hipSuccess); } } hipError_t hipDeviceSynchronize(void) { HIP_INIT_SPECIAL_API(hipDeviceSynchronize, TRACE_SYNC); - return ihipLogStatus(ihipSynchronize()); + return ihipLogStatus(ihipSynchronize(tls)); } hipError_t hipDeviceReset(void) { @@ -171,7 +171,7 @@ hipError_t hipDeviceReset(void) { } -hipError_t ihipDeviceSetState(void) { +hipError_t ihipDeviceSetState(TlsData *tls) { hipError_t e = hipErrorInvalidContext; auto* ctx = ihipGetTlsDefaultCtx(); @@ -305,6 +305,12 @@ hipError_t ihipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device *hdp = prop->hdpRegFlushCntl; } break; + case hipDeviceAttributeCooperativeLaunch: + *pi = prop->cooperativeLaunch; + break; + case hipDeviceAttributeCooperativeMultiDeviceLaunch: + *pi = prop->cooperativeMultiDeviceLaunch; + break; default: e = hipErrorInvalidValue; break; diff --git a/src/hip_error.cpp b/src/hip_error.cpp index 6f1184d92f..0f876b4f26 100644 --- a/src/hip_error.cpp +++ b/src/hip_error.cpp @@ -33,8 +33,8 @@ hipError_t hipGetLastError() { HIP_INIT_API(hipGetLastError); // Return last error, but then reset the state: - hipError_t e = ihipLogStatus(tls_lastHipError); - tls_lastHipError = hipSuccess; + hipError_t e = ihipLogStatus(tls->lastHipError); + tls->lastHipError = hipSuccess; return e; } @@ -42,7 +42,7 @@ hipError_t hipPeekAtLastError() { HIP_INIT_API(hipPeekAtLastError); // peek at last error, but don't reset it. - return ihipLogStatus(tls_lastHipError); + return ihipLogStatus(tls->lastHipError); } const char* hipGetErrorName(hipError_t hip_error) { diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index f3d6eae6b7..dcfdfb6137 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -139,24 +139,12 @@ std::atomic g_lastShortTid(1); std::vector g_dbStartTriggers; std::vector g_dbStopTriggers; -//================================================================================================= -// Thread-local storage: -//================================================================================================= - -// This is the implicit context used by all HIP commands. -// It can be set by hipSetDevice or by the CTX manipulation commands: -thread_local hipError_t tls_lastHipError = hipSuccess; - - -thread_local TidInfo tls_tidInfo; - - //================================================================================================= // Top-level "free" functions: //================================================================================================= -uint64_t recordApiTrace(std::string* fullStr, const std::string& apiStr) { - auto apiSeqNum = tls_tidInfo.apiSeqNum(); - auto tid = tls_tidInfo.tid(); +uint64_t recordApiTrace(TlsData *tls, std::string* fullStr, const std::string& apiStr) { + auto apiSeqNum = tls->tidInfo.apiSeqNum(); + auto tid = tls->tidInfo.tid(); if ((tid < g_dbStartTriggers.size()) && (apiSeqNum >= g_dbStartTriggers[tid].nextTrigger())) { printf("info: resume profiling at %lu\n", apiSeqNum); @@ -179,7 +167,7 @@ uint64_t recordApiTrace(std::string* fullStr, const std::string& apiStr) { if (COMPILE_HIP_DB && HIP_TRACE_API) { - fprintf(stderr, "%s<c_str(), apiStartTick, + fprintf(stderr, "%s<tidInfo.pid(), fullStr->c_str(), apiStartTick, API_COLOR_END); } @@ -206,32 +194,18 @@ ihipCtx_t* ihipGetPrimaryCtx(unsigned deviceIndex) { return device ? device->getPrimaryCtx() : NULL; }; - -static thread_local ihipCtx_t* tls_defaultCtx = nullptr; -void ihipSetTlsDefaultCtx(ihipCtx_t* ctx) { tls_defaultCtx = ctx; } - - -//--- -// TODO - review the context creation strategy here. Really should be: -// - first "non-device" runtime call creates the context for this thread. Allowed to call -// setDevice first. -// - hipDeviceReset destroys the primary context for device? -// - Then context is created again for next usage. -ihipCtx_t* ihipGetTlsDefaultCtx() { - // Per-thread initialization of the TLS: - if ((tls_defaultCtx == nullptr) && (g_deviceCnt > 0)) { - ihipSetTlsDefaultCtx(ihipGetPrimaryCtx(0)); - } - return tls_defaultCtx; -} - -hipError_t ihipSynchronize(void) { +hipError_t ihipSynchronize(TlsData *tls) { ihipGetTlsDefaultCtx()->locked_waitAllStreams(); // ignores non-blocking streams, this waits // for all activity to finish. return (hipSuccess); } +TlsData* tls_get_ptr() { + static thread_local TlsData data; + return &data; +} + //================================================================================================= // ihipStream_t: //================================================================================================= @@ -904,6 +878,10 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) { prop->integrated = 1; } + // Enable the cooperative group for gfx9+ + prop->cooperativeLaunch = (prop->gcnArch < 900) ? 0 : 1; + prop->cooperativeMultiDeviceLaunch = (prop->gcnArch < 900) ? 0 : 1; + err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_EXT_AGENT_INFO_IMAGE_1D_MAX_ELEMENTS, &prop->maxTexture1D); DeviceErrorCheck(err); @@ -1478,7 +1456,7 @@ hipError_t hip_init() { } } -hipError_t ihipStreamSynchronize(hipStream_t stream) { +hipError_t ihipStreamSynchronize(TlsData *tls, hipStream_t stream) { hipError_t e = hipSuccess; if (stream == hipStreamNull) { @@ -1499,7 +1477,8 @@ void ihipStreamCallbackHandler(ihipStreamCallback_t* cb) { // Synchronize stream tprintf(DB_SYNC, "ihipStreamCallbackHandler wait on stream %s\n", ToString(cb->_stream).c_str()); - e = ihipStreamSynchronize(cb->_stream); + GET_TLS(); + e = ihipStreamSynchronize(tls, cb->_stream); // Call registered callback function cb->_callback(cb->_stream, e, cb->_userData); @@ -1514,6 +1493,7 @@ void ihipStreamCallbackHandler(ihipStreamCallback_t* cb) { hipStream_t ihipSyncAndResolveStream(hipStream_t stream, bool lockAcquired) { if (stream == hipStreamNull) { // Submitting to NULL stream, call locked_syncDefaultStream to wait for all other streams: + GET_TLS(); ihipCtx_t* ctx = ihipGetTlsDefaultCtx(); tprintf(DB_SYNC, "ihipSyncAndResolveStream %s wait on default stream\n", ToString(stream).c_str()); @@ -1577,15 +1557,16 @@ void ihipPrintKernelLaunch(const char* kernelName, const grid_launch_parm* lp, const hipStream_t stream) { if ((HIP_TRACE_API & (1 << TRACE_KCMD)) || HIP_PROFILE_API || (COMPILE_HIP_DB & HIP_TRACE_API)) { + GET_TLS(); std::stringstream os; - os << tls_tidInfo.pid() << " " << tls_tidInfo.tid() << "." << tls_tidInfo.apiSeqNum() << " hipLaunchKernel '" + os << tls->tidInfo.pid() << " " << tls->tidInfo.tid() << "." << tls->tidInfo.apiSeqNum() << " hipLaunchKernel '" << kernelName << "'" << " gridDim:" << lp->grid_dim << " groupDim:" << lp->group_dim << " sharedMem:+" << lp->dynamic_group_mem_bytes << " " << *stream; if (COMPILE_HIP_DB && HIP_TRACE_API) { std::string fullStr; - recordApiTrace(&fullStr, os.str()); + recordApiTrace(tls, &fullStr, os.str()); } if (HIP_PROFILE_API == 0x1) { @@ -2513,6 +2494,8 @@ namespace hip_impl { [[noreturn]] void hip_throw(const std::exception& ex) { #if defined(__cpp_exceptions) + if (auto rte = dynamic_cast(&ex)) throw *rte; + if (auto lge = dynamic_cast(&ex)) throw *lge; throw ex; #else std::cerr << ex.what() << std::endl; diff --git a/src/hip_hcc_internal.h b/src/hip_hcc_internal.h index 7ab3fa6de1..d276aa9caa 100644 --- a/src/hip_hcc_internal.h +++ b/src/hip_hcc_internal.h @@ -130,10 +130,26 @@ struct ProfTrigger { //--- -// Extern tls -extern thread_local hipError_t tls_lastHipError; -extern thread_local TidInfo tls_tidInfo; -extern thread_local bool tls_getPrimaryCtx; +// Extern TLS +// Use a single struct to hold all TLS data. Attempt to reduce TLS accesses. +struct TlsData { + explicit TlsData() { + lastHipError = hipSuccess; + getPrimaryCtx = true; + defaultCtx = nullptr; + } + + hipError_t lastHipError; + TidInfo tidInfo; + // This is the implicit context used by all HIP commands. + // It can be set by hipSetDevice or by the CTX manipulation commands: + ihipCtx_t* defaultCtx; + // Stack of contexts + std::stack ctxStack; + bool getPrimaryCtx; +}; +TlsData* tls_get_ptr(); +#define GET_TLS() TlsData *tls = tls_get_ptr() extern std::vector g_dbStartTriggers; extern std::vector g_dbStopTriggers; @@ -194,7 +210,7 @@ extern const char* API_COLOR_END; // Compile support for trace markers that are displayed on CodeXL GUI at start/stop of each function // boundary. // TODO - currently we print the trace message at the beginning. if we waited, we could also -// tls_tidInfo return codes, and any values returned through ptr-to-args (ie the pointers allocated +// tls->tidInfo return codes, and any values returned through ptr-to-args (ie the pointers allocated // by hipMalloc). #if COMPILE_HIP_ATP_MARKER #include "CXLActivityLogger.h" @@ -251,10 +267,11 @@ static const DbName dbName[] = { #define tprintf(trace_level, ...) \ { \ if (HIP_DB & (1 << (trace_level))) { \ + GET_TLS(); \ char msgStr[1000]; \ snprintf(msgStr, sizeof(msgStr), __VA_ARGS__); \ fprintf(stderr, " %ship-%s pid:%d tid:%d:%s%s", dbName[trace_level]._color, \ - dbName[trace_level]._shortName, tls_tidInfo.pid(), tls_tidInfo.tid(), msgStr, KNRM); \ + dbName[trace_level]._shortName, tls->tidInfo.pid(), tls->tidInfo.tid(), msgStr, KNRM); \ } \ } #else @@ -266,18 +283,19 @@ static const DbName dbName[] = { static inline uint64_t getTicks() { return hc::get_system_ticks(); } //--- -extern uint64_t recordApiTrace(std::string* fullStr, const std::string& apiStr); +extern uint64_t recordApiTrace(TlsData *tls, std::string* fullStr, const std::string& apiStr); #if COMPILE_HIP_ATP_MARKER || (COMPILE_HIP_TRACE_API & 0x1) #define API_TRACE(forceTrace, ...) \ + GET_TLS(); \ uint64_t hipApiStartTick = 0; \ { \ - tls_tidInfo.incApiSeqNum(); \ + tls->tidInfo.incApiSeqNum(); \ if (forceTrace || \ (HIP_PROFILE_API || (COMPILE_HIP_DB && (HIP_TRACE_API & (1 << TRACE_ALL))))) { \ std::string apiStr = std::string(__func__) + " (" + ToString(__VA_ARGS__) + ')'; \ std::string fullStr; \ - hipApiStartTick = recordApiTrace(&fullStr, apiStr); \ + hipApiStartTick = recordApiTrace(tls, &fullStr, apiStr); \ if (HIP_PROFILE_API == 0x1) { \ MARKER_BEGIN(__func__, "HIP") \ } else if (HIP_PROFILE_API == 0x2) { \ @@ -288,10 +306,13 @@ extern uint64_t recordApiTrace(std::string* fullStr, const std::string& apiStr); #else // Swallow API_TRACE -#define API_TRACE(IS_CMD, ...) tls_tidInfo.incApiSeqNum(); +#define API_TRACE(IS_CMD, ...) GET_TLS(); tls->tidInfo.incApiSeqNum(); #endif -#define HIP_SET_DEVICE() ihipDeviceSetState(); +#define ihipGetTlsDefaultCtx() iihipGetTlsDefaultCtx(tls) +#define ihipSetTlsDefaultCtx(ctx) tls->defaultCtx = ctx + +#define HIP_SET_DEVICE() ihipDeviceSetState(tls); // This macro should be called at the beginning of every HIP API. // It initializes the hip runtime (exactly once), and @@ -314,22 +335,22 @@ extern uint64_t recordApiTrace(std::string* fullStr, const std::string& apiStr); // This macro should be called at the end of every HIP API, and only at the end of top-level hip // APIS (not internal hip) It has dual function: logs the last error returned for use by // hipGetLastError, and also prints the closing message when the debug trace is enabled. -#define ihipLogStatus(hipStatus) \ - ({ \ - hipError_t localHipStatus = hipStatus; /*local copy so hipStatus only evaluated once*/ \ - tls_lastHipError = localHipStatus; \ - \ - if ((COMPILE_HIP_TRACE_API & 0x2) && HIP_TRACE_API & (1 << TRACE_ALL)) { \ - auto ticks = getTicks() - hipApiStartTick; \ - fprintf(stderr, " %ship-api pid:%d tid:%d.%lu %-30s ret=%2d (%s)>> +%lu ns%s\n", \ - (localHipStatus == 0) ? API_COLOR : KRED, tls_tidInfo.pid(), tls_tidInfo.tid(), \ - tls_tidInfo.apiSeqNum(), __func__, localHipStatus, \ - ihipErrorString(localHipStatus), ticks, API_COLOR_END); \ - } \ - if (HIP_PROFILE_API) { \ - MARKER_END(); \ - } \ - localHipStatus; \ +#define ihipLogStatus(hipStatus) \ + ({ \ + hipError_t localHipStatus = hipStatus; /*local copy so hipStatus only evaluated once*/ \ + tls->lastHipError = localHipStatus; \ + \ + if ((COMPILE_HIP_TRACE_API & 0x2) && HIP_TRACE_API & (1 << TRACE_ALL)) { \ + auto ticks = getTicks() - hipApiStartTick; \ + fprintf(stderr, " %ship-api pid:%d tid:%d.%lu %-30s ret=%2d (%s)>> +%lu ns%s\n", \ + (localHipStatus == 0) ? API_COLOR : KRED, tls->tidInfo.pid(), tls->tidInfo.tid(), \ + tls->tidInfo.apiSeqNum(), __func__, localHipStatus, \ + ihipErrorString(localHipStatus), ticks, API_COLOR_END); \ + } \ + if (HIP_PROFILE_API) { \ + MARKER_END(); \ + } \ + localHipStatus; \ }) @@ -936,11 +957,9 @@ extern hsa_agent_t* g_allAgents; // CPU agents + all the visible GPU agents. // Extern functions: extern void ihipInit(); extern const char* ihipErrorString(hipError_t); -extern ihipCtx_t* ihipGetTlsDefaultCtx(); -extern void ihipSetTlsDefaultCtx(ihipCtx_t* ctx); -extern hipError_t ihipSynchronize(void); +extern hipError_t ihipSynchronize(TlsData *tls); extern void ihipCtxStackUpdate(); -extern hipError_t ihipDeviceSetState(); +extern hipError_t ihipDeviceSetState(TlsData *tls); extern ihipDevice_t* ihipGetDevice(int); ihipCtx_t* ihipGetPrimaryCtx(unsigned deviceIndex); @@ -949,7 +968,7 @@ hipError_t hipModuleGetFunctionEx(hipFunction_t* hfunc, hipModule_t hmod, hipStream_t ihipSyncAndResolveStream(hipStream_t, bool lockAcquired = 0); -hipError_t ihipStreamSynchronize(hipStream_t stream); +hipError_t ihipStreamSynchronize(TlsData *tls, hipStream_t stream); void ihipStreamCallbackHandler(ihipStreamCallback_t* cb); // Stream printf functions: @@ -1002,5 +1021,18 @@ hipError_t memcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKi hipStream_t stream); }; +//--- +// TODO - review the context creation strategy here. Really should be: +// - first "non-device" runtime call creates the context for this thread. Allowed to call +// setDevice first. +// - hipDeviceReset destroys the primary context for device? +// - Then context is created again for next usage. +static inline ihipCtx_t* iihipGetTlsDefaultCtx(TlsData* tls) { + // Per-thread initialization of the TLS: + if ((tls->defaultCtx == nullptr) && (g_deviceCnt > 0)) { + tls->defaultCtx = ihipGetPrimaryCtx(0); + } + return tls->defaultCtx; +} #endif diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index 99ea4b5b24..61f34db244 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -300,7 +300,7 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag return ihipLogStatus(hip_status); } -hipError_t ihipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { +hipError_t ihipHostMalloc(TlsData *tls, void** ptr, size_t sizeBytes, unsigned int flags) { hipError_t hip_status = hipSuccess; if (HIP_SYNC_HOST_ALLOC) { @@ -372,7 +372,7 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) { HIP_INIT_SPECIAL_API(hipHostMalloc, (TRACE_MEM), ptr, sizeBytes, flags); HIP_SET_DEVICE(); hipError_t hip_status = hipSuccess; - hip_status = ihipHostMalloc(ptr, sizeBytes, flags); + hip_status = ihipHostMalloc(tls, ptr, sizeBytes, flags); return ihipLogStatus(hip_status); } @@ -383,7 +383,7 @@ hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags) { if(flags != hipMemAttachGlobal) hip_status = hipErrorInvalidValue; else - hip_status = ihipHostMalloc(devPtr, size, hipHostMallocDefault); + hip_status = ihipHostMalloc(tls, devPtr, size, hipHostMallocDefault); return ihipLogStatus(hip_status); } @@ -397,7 +397,7 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) { }; // width in bytes -hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height, size_t depth) { +hipError_t ihipMallocPitch(TlsData* tls, void** ptr, size_t* pitch, size_t width, size_t height, size_t depth) { hipError_t hip_status = hipSuccess; if(ptr==NULL || pitch == NULL) { @@ -460,7 +460,7 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height if (width == 0 || height == 0) return ihipLogStatus(hipErrorUnknown); - hip_status = ihipMallocPitch(ptr, pitch, width, height, 0); + hip_status = ihipMallocPitch(tls, ptr, pitch, width, height, 0); return ihipLogStatus(hip_status); } @@ -475,7 +475,7 @@ hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) { size_t pitch; hip_status = - ihipMallocPitch(&pitchedDevPtr->ptr, &pitch, extent.width, extent.height, extent.depth); + ihipMallocPitch(tls, &pitchedDevPtr->ptr, &pitch, extent.width, extent.height, extent.depth); if (hip_status == hipSuccess) { pitchedDevPtr->pitch = pitch; pitchedDevPtr->xsize = extent.width; diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 9e972da246..04cdedbb75 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -137,7 +137,7 @@ hipError_t hipModuleUnload(hipModule_t hmod) { // TODO - improve this synchronization so it is thread-safe. // Currently we want for all inflight activity to complete, but don't prevent another // thread from launching new kernels before we finish this operation. - ihipSynchronize(); + ihipSynchronize(tls); delete hmod; // The ihipModule_t dtor will clean everything up. hmod = nullptr; @@ -145,7 +145,7 @@ hipError_t hipModuleUnload(hipModule_t hmod) { return ihipLogStatus(hipSuccess); } -hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, +hipError_t ihipModuleLaunchKernel(TlsData *tls, hipFunction_t f, uint32_t globalWorkSizeX, uint32_t globalWorkSizeY, uint32_t globalWorkSizeZ, uint32_t localWorkSizeX, uint32_t localWorkSizeY, uint32_t localWorkSizeZ, size_t sharedMemBytes, @@ -285,7 +285,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, uint32_t gridDimX, uint32_t gr void** kernelParams, void** extra) { HIP_INIT_API(hipModuleLaunchKernel, f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra); - return ihipLogStatus(ihipModuleLaunchKernel( + return ihipLogStatus(ihipModuleLaunchKernel(tls, f, blockDimX * gridDimX, blockDimY * gridDimY, gridDimZ * blockDimZ, blockDimX, blockDimY, blockDimZ, sharedMemBytes, hStream, kernelParams, extra, nullptr, nullptr, 0)); } @@ -298,7 +298,7 @@ hipError_t hipExtModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags) { HIP_INIT_API(hipExtModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra); - return ihipLogStatus(ihipModuleLaunchKernel( + return ihipLogStatus(ihipModuleLaunchKernel(tls, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, flags)); } @@ -311,14 +311,14 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, hipEvent_t startEvent, hipEvent_t stopEvent) { HIP_INIT_API(hipHccModuleLaunchKernel, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra); - return ihipLogStatus(ihipModuleLaunchKernel( + return ihipLogStatus(ihipModuleLaunchKernel(tls, f, globalWorkSizeX, globalWorkSizeY, globalWorkSizeZ, localWorkSizeX, localWorkSizeY, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, 0)); } hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices, unsigned int flags) { - HIP_INIT_API(hipExtLaunchMultiKernelMultiDevice, launchParamsList,numDevices,flags); + HIP_INIT_API(hipExtLaunchMultiKernelMultiDevice, launchParamsList, numDevices, flags); hipError_t result; if ((numDevices > g_deviceCnt) || (launchParamsList == nullptr)) { @@ -361,7 +361,7 @@ hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, for (int i = 0; i < numDevices; ++i) { const hipLaunchParams& lp = launchParamsList[i]; - result = ihipModuleLaunchKernel(kds[i], + result = ihipModuleLaunchKernel(tls, kds[i], lp.gridDim.x * lp.blockDim.x, lp.gridDim.y * lp.blockDim.y, lp.gridDim.z * lp.blockDim.z, @@ -395,6 +395,7 @@ namespace hip_impl { } hsa_agent_t this_agent() { + GET_TLS(); auto ctx = ihipGetTlsDefaultCtx(); if (!ctx) throw runtime_error{"No active HIP context."}; @@ -574,6 +575,7 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, namespace { inline void track(const hip_impl::Agent_global& x, hsa_agent_t agent) { + GET_TLS(); tprintf(DB_MEM, " add variable '%s' with ptr=%p size=%u to tracker\n", x.name, x.address, x.byte_cnt); @@ -693,7 +695,7 @@ namespace hip_impl { } } // Namespace hip_impl. -hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const char* name, +hipError_t ihipModuleGetFunction(TlsData *tls, hipFunction_t* func, hipModule_t hmod, const char* name, hsa_agent_t *agent = nullptr) { using namespace hip_impl; @@ -728,14 +730,14 @@ hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const ch // Get kernel for the current hsa agent. hipError_t hipModuleGetFunction(hipFunction_t* hfunc, hipModule_t hmod, const char* name) { HIP_INIT_API(hipModuleGetFunction, hfunc, hmod, name); - return ihipLogStatus(ihipModuleGetFunction(hfunc, hmod, name)); + return ihipLogStatus(ihipModuleGetFunction(tls, hfunc, hmod, name)); } // Get kernel for the given hsa agent. Internal use only. hipError_t hipModuleGetFunctionEx(hipFunction_t* hfunc, hipModule_t hmod, const char* name, hsa_agent_t *agent) { HIP_INIT_API(hipModuleGetFunctionEx, hfunc, hmod, name, agent); - return ihipLogStatus(ihipModuleGetFunction(hfunc, hmod, name, agent)); + return ihipLogStatus(ihipModuleGetFunction(tls, hfunc, hmod, name, agent)); } namespace { @@ -743,7 +745,7 @@ const amd_kernel_code_v3_t *header_v3(const ihipModuleSymbol_t& kd) { return reinterpret_cast(kd._header); } -hipFuncAttributes make_function_attributes(const ihipModuleSymbol_t& kd) { +hipFuncAttributes make_function_attributes(TlsData *tls, const ihipModuleSymbol_t& kd) { hipFuncAttributes r{}; hipDeviceProp_t prop{}; @@ -788,20 +790,20 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func) HIP_INIT_API(hipFuncGetAttributes, attr, func); using namespace hip_impl; - if (!attr) return hipErrorInvalidValue; - if (!func) return hipErrorInvalidDeviceFunction; + if (!attr) return ihipLogStatus(hipErrorInvalidValue); + if (!func) return ihipLogStatus(hipErrorInvalidDeviceFunction); auto agent = this_agent(); auto kd = get_program_state().kernel_descriptor(reinterpret_cast(func), agent); if (!kd->_header) throw runtime_error{"Ill-formed Kernel_descriptor."}; - *attr = make_function_attributes(*kd); + *attr = make_function_attributes(tls, *kd); return ihipLogStatus(hipSuccess); } -hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) { +hipError_t ihipModuleLoadData(TlsData *tls, hipModule_t* module, const void* image) { using namespace hip_impl; if (!module) return hipErrorInvalidValue; @@ -839,7 +841,7 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) { hipError_t hipModuleLoadData(hipModule_t* module, const void* image) { HIP_INIT_API(hipModuleLoadData, module, image); - return ihipLogStatus(ihipModuleLoadData(module,image)); + return ihipLogStatus(ihipModuleLoadData(tls,module,image)); } hipError_t hipModuleLoad(hipModule_t* module, const char* fname) { @@ -853,13 +855,13 @@ hipError_t hipModuleLoad(hipModule_t* module, const char* fname) { vector tmp{istreambuf_iterator{file}, istreambuf_iterator{}}; - return ihipLogStatus(ihipModuleLoadData(module, tmp.data())); + return ihipLogStatus(ihipModuleLoadData(tls, module, tmp.data())); } hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* image, unsigned int numOptions, hipJitOption* options, void** optionValues) { HIP_INIT_API(hipModuleLoadDataEx, module, image, numOptions, options, optionValues); - return ihipLogStatus(ihipModuleLoadData(module, image)); + return ihipLogStatus(ihipModuleLoadData(tls, module, image)); } hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name) { @@ -903,7 +905,7 @@ void getGprsLdsUsage(hipFunction_t f, size_t* usedVGPRS, size_t* usedSGPRS, size } } -hipError_t ihipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* blockSize, +hipError_t ihipOccupancyMaxPotentialBlockSize(TlsData *tls, uint32_t* gridSize, uint32_t* blockSize, hipFunction_t f, size_t dynSharedMemPerBlk, uint32_t blockSizeLimit) { @@ -1021,12 +1023,12 @@ hipError_t hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* block { HIP_INIT_API(hipOccupancyMaxPotentialBlockSize, gridSize, blockSize, f, dynSharedMemPerBlk, blockSizeLimit); - return ihipLogStatus(ihipOccupancyMaxPotentialBlockSize( + return ihipLogStatus(ihipOccupancyMaxPotentialBlockSize(tls, gridSize, blockSize, f, dynSharedMemPerBlk, blockSizeLimit)); } hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor( - uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk) + TlsData *tls, uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk) { using namespace hip_impl; @@ -1083,7 +1085,7 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor( HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessor, numBlocks, f, blockSize, dynSharedMemPerBlk); return ihipLogStatus(ihipOccupancyMaxActiveBlocksPerMultiprocessor( - numBlocks, f, blockSize, dynSharedMemPerBlk)); + tls, numBlocks, f, blockSize, dynSharedMemPerBlk)); } hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( @@ -1093,5 +1095,5 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags, numBlocks, f, blockSize, dynSharedMemPerBlk, flags); return ihipLogStatus(ihipOccupancyMaxActiveBlocksPerMultiprocessor( - numBlocks, f, blockSize, dynSharedMemPerBlk)); + tls, numBlocks, f, blockSize, dynSharedMemPerBlk)); } diff --git a/src/hip_peer.cpp b/src/hip_peer.cpp index f2fca94404..7781af1dbe 100644 --- a/src/hip_peer.cpp +++ b/src/hip_peer.cpp @@ -84,7 +84,7 @@ hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t //--- // Disable visibility of this device into memory allocated on peer device. // Remove this device from peer device peerlist. -hipError_t ihipDisablePeerAccess(hipCtx_t peerCtx) { +hipError_t ihipDisablePeerAccess(TlsData* tls, hipCtx_t peerCtx) { hipError_t err = hipSuccess; auto thisCtx = ihipGetTlsDefaultCtx(); @@ -119,7 +119,7 @@ hipError_t ihipDisablePeerAccess(hipCtx_t peerCtx) { //--- // Allow the current device to see all memory allocated on peerCtx. // This should add this device to the peer-device peer list. -hipError_t ihipEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) { +hipError_t ihipEnablePeerAccess(TlsData* tls, hipCtx_t peerCtx, unsigned int flags) { hipError_t err = hipSuccess; if (flags != 0) { err = hipErrorInvalidValue; @@ -186,14 +186,14 @@ hipError_t hipDeviceCanAccessPeer(int* canAccessPeer, int deviceId, int peerDevi hipError_t hipDeviceDisablePeerAccess(int peerDeviceId) { HIP_INIT_API(hipDeviceDisablePeerAccess, peerDeviceId); - return ihipLogStatus(ihipDisablePeerAccess(ihipGetPrimaryCtx(peerDeviceId))); + return ihipLogStatus(ihipDisablePeerAccess(tls, ihipGetPrimaryCtx(peerDeviceId))); } hipError_t hipDeviceEnablePeerAccess(int peerDeviceId, unsigned int flags) { HIP_INIT_API(hipDeviceEnablePeerAccess, peerDeviceId, flags); - return ihipLogStatus(ihipEnablePeerAccess(ihipGetPrimaryCtx(peerDeviceId), flags)); + return ihipLogStatus(ihipEnablePeerAccess(tls, ihipGetPrimaryCtx(peerDeviceId), flags)); } @@ -214,11 +214,11 @@ hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int src hipError_t hipCtxEnablePeerAccess(hipCtx_t peerCtx, unsigned int flags) { HIP_INIT_API(hipCtxEnablePeerAccess, peerCtx, flags); - return ihipLogStatus(ihipEnablePeerAccess(peerCtx, flags)); + return ihipLogStatus(ihipEnablePeerAccess(tls, peerCtx, flags)); } hipError_t hipCtxDisablePeerAccess(hipCtx_t peerCtx) { HIP_INIT_API(hipCtxDisablePeerAccess, peerCtx); - return ihipLogStatus(ihipDisablePeerAccess(peerCtx)); + return ihipLogStatus(ihipDisablePeerAccess(tls, peerCtx)); } diff --git a/src/hip_stream.cpp b/src/hip_stream.cpp index 7812530dda..e3e4975b7e 100644 --- a/src/hip_stream.cpp +++ b/src/hip_stream.cpp @@ -48,7 +48,7 @@ enum queue_priority #endif //--- -hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priority) { +hipError_t ihipStreamCreate(TlsData *tls, hipStream_t* stream, unsigned int flags, int priority) { ihipCtx_t* ctx = ihipGetTlsDefaultCtx(); hipError_t e = hipSuccess; @@ -97,7 +97,7 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags) { HIP_INIT_API(hipStreamCreateWithFlags, stream, flags); if(flags == hipStreamDefault || flags == hipStreamNonBlocking) - return ihipLogStatus(ihipStreamCreate(stream, flags, priority_normal)); + return ihipLogStatus(ihipStreamCreate(tls, stream, flags, priority_normal)); else return ihipLogStatus(hipErrorInvalidValue); } @@ -106,7 +106,7 @@ hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags) { hipError_t hipStreamCreate(hipStream_t* stream) { HIP_INIT_API(hipStreamCreate, stream); - return ihipLogStatus(ihipStreamCreate(stream, hipStreamDefault, priority_normal)); + return ihipLogStatus(ihipStreamCreate(tls, stream, hipStreamDefault, priority_normal)); } //--- @@ -115,7 +115,7 @@ hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, // clamp priority to range [priority_high:priority_low] priority = (priority < priority_high ? priority_high : (priority > priority_low ? priority_low : priority)); - return ihipLogStatus(ihipStreamCreate(stream, flags, priority)); + return ihipLogStatus(ihipStreamCreate(tls, stream, flags, priority)); } //--- @@ -183,7 +183,7 @@ hipError_t hipStreamQuery(hipStream_t stream) { hipError_t hipStreamSynchronize(hipStream_t stream) { HIP_INIT_SPECIAL_API(hipStreamSynchronize, TRACE_SYNC, stream); - return ihipLogStatus(ihipStreamSynchronize(stream)); + return ihipLogStatus(ihipStreamSynchronize(tls, stream)); } diff --git a/src/hip_texture.cpp b/src/hip_texture.cpp index e6cd352008..fe97895381 100644 --- a/src/hip_texture.cpp +++ b/src/hip_texture.cpp @@ -381,9 +381,10 @@ hipError_t hipGetTextureObjectTextureDesc(hipTextureDesc* pTexDesc, } // Texture Reference APIs -hipError_t ihipBindTextureImpl(int dim, enum hipTextureReadMode readMode, size_t* offset, +hipError_t ihipBindTextureImpl(TlsData *tls_, int dim, enum hipTextureReadMode readMode, size_t* offset, const void* devPtr, const struct hipChannelFormatDesc* desc, size_t size, textureReference* tex) { + TlsData *tls = (tls_ == nullptr) ? tls_get_ptr() : tls_; hipError_t hip_status = hipSuccess; enum hipTextureAddressMode addressMode = tex->addressMode[0]; enum hipTextureFilterMode filterMode = tex->filterMode; @@ -447,12 +448,12 @@ hipError_t hipBindTexture(size_t* offset, textureReference* tex, const void* dev HIP_INIT_API(hipBindTexture, offset, tex, devPtr, desc, size); hipError_t hip_status = hipSuccess; // TODO: hipReadModeElementType is default. - hip_status = ihipBindTextureImpl(hipTextureType1D, hipReadModeElementType, offset, devPtr, desc, + hip_status = ihipBindTextureImpl(tls, hipTextureType1D, hipReadModeElementType, offset, devPtr, desc, size, tex); return ihipLogStatus(hip_status); } -hipError_t ihipBindTexture2DImpl(int dim, enum hipTextureReadMode readMode, size_t* offset, +hipError_t ihipBindTexture2DImpl(TlsData *tls, int dim, enum hipTextureReadMode readMode, size_t* offset, const void* devPtr, const struct hipChannelFormatDesc* desc, size_t width, size_t height, textureReference* tex) { hipError_t hip_status = hipSuccess; @@ -519,15 +520,16 @@ hipError_t hipBindTexture2D(size_t* offset, textureReference* tex, const void* d size_t pitch) { HIP_INIT_API(hipBindTexture2D, offset, tex, devPtr, desc, width, height, pitch); hipError_t hip_status = hipSuccess; - hip_status = ihipBindTexture2DImpl(hipTextureType2D, hipReadModeElementType, offset, devPtr, + hip_status = ihipBindTexture2DImpl(tls, hipTextureType2D, hipReadModeElementType, offset, devPtr, desc, width, height, tex); return ihipLogStatus(hip_status); } -hipError_t ihipBindTextureToArrayImpl(int dim, enum hipTextureReadMode readMode, +hipError_t ihipBindTextureToArrayImpl(TlsData *tls_, int dim, enum hipTextureReadMode readMode, hipArray_const_t array, const struct hipChannelFormatDesc& desc, textureReference* tex) { + TlsData *tls = (tls_ == nullptr) ? tls_get_ptr() : tls_; hipError_t hip_status = hipSuccess; enum hipTextureAddressMode addressMode = tex->addressMode[0]; enum hipTextureFilterMode filterMode = tex->filterMode; @@ -617,7 +619,7 @@ hipError_t hipBindTextureToArray(textureReference* tex, hipArray_const_t array, hipError_t hip_status = hipSuccess; // TODO: hipReadModeElementType is default. hip_status = - ihipBindTextureToArrayImpl(array->textureType, hipReadModeElementType, array, *desc, tex); + ihipBindTextureToArrayImpl(tls, array->textureType, hipReadModeElementType, array, *desc, tex); return ihipLogStatus(hip_status); } @@ -629,7 +631,7 @@ hipError_t hipBindTextureToMipmappedArray(textureReference* tex, return ihipLogStatus(hip_status); } -hipError_t ihipUnbindTextureImpl(const hipTextureObject_t& textureObject) { +hipError_t ihipUnbindTextureImpl(TlsData *tls, const hipTextureObject_t& textureObject) { hipError_t hip_status = hipSuccess; auto ctx = ihipGetTlsDefaultCtx(); @@ -654,7 +656,7 @@ hipError_t ihipUnbindTextureImpl(const hipTextureObject_t& textureObject) { hipError_t hipUnbindTexture(const textureReference* tex) { HIP_INIT_API(hipUnbindTexture, tex); hipError_t hip_status = hipSuccess; - hip_status = ihipUnbindTextureImpl(tex->textureObject); + hip_status = ihipUnbindTextureImpl(tls, tex->textureObject); return ihipLogStatus(hip_status); } @@ -726,7 +728,7 @@ hipError_t hipTexRefSetArray(textureReference* tex, hipArray_const_t array, unsi HIP_INIT_API(hipTexRefSetArray, tex, array, flags); hipError_t hip_status = hipSuccess; - hip_status = ihipBindTextureToArrayImpl(array->textureType, hipReadModeElementType, array, + hip_status = ihipBindTextureToArrayImpl(tls, array->textureType, hipReadModeElementType, array, array->desc, tex); return ihipLogStatus(hip_status); } @@ -737,7 +739,7 @@ hipError_t hipTexRefSetAddress(size_t* offset, textureReference* tex, hipDevicep HIP_INIT_API(hipTexRefSetAddress, offset, tex, devPtr, size); hipError_t hip_status = hipSuccess; // TODO: hipReadModeElementType is default. - hip_status = ihipBindTextureImpl(hipTextureType1D, hipReadModeElementType, offset, devPtr, NULL, + hip_status = ihipBindTextureImpl(tls, hipTextureType1D, hipReadModeElementType, offset, devPtr, NULL, size, tex); return ihipLogStatus(hip_status); } @@ -748,7 +750,7 @@ hipError_t hipTexRefSetAddress2D(textureReference* tex, const HIP_ARRAY_DESCRIPT size_t offset; hipError_t hip_status = hipSuccess; // TODO: hipReadModeElementType is default. - hip_status = ihipBindTexture2DImpl(hipTextureType2D, hipReadModeElementType, &offset, devPtr, + hip_status = ihipBindTexture2DImpl(tls, hipTextureType2D, hipReadModeElementType, &offset, devPtr, NULL, desc->Width, desc->Height, tex); return ihipLogStatus(hip_status); } diff --git a/src/macro_based_grid_launch.inl b/src/macro_based_grid_launch.inl index 1e36903c56..2e804f090b 100644 --- a/src/macro_based_grid_launch.inl +++ b/src/macro_based_grid_launch.inl @@ -71,7 +71,8 @@ namespace hip_impl if (COMPILE_HIP_DB && HIP_TRACE_API) { std::string fullStr; - recordApiTrace(&fullStr, os.str()); + GET_TLS(); + recordApiTrace(tls, &fullStr, os.str()); } } } diff --git a/tests/hipify-clang/unit_tests/headers/headers_test_10.cu b/tests/hipify-clang/unit_tests/headers/headers_test_10.cu new file mode 100644 index 0000000000..1c2db50e3b --- /dev/null +++ b/tests/hipify-clang/unit_tests/headers/headers_test_10.cu @@ -0,0 +1,14 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args + +// Checks that HIP header file is included after #pragma once, +// which goes before include guard controlling macro. +// CHECK: #pragma once +// CHECK-NEXT: #include +#pragma once +#ifndef HEADERS_TEST_10_H +// CHECK: #ifndef HEADERS_TEST_10_H +// CHECK-NOT: #include +#define HEADERS_TEST_10_H +#include +static int counter = 0; +#endif // HEADERS_TEST_10_H diff --git a/tests/hipify-clang/unit_tests/headers/headers_test_11.cu b/tests/hipify-clang/unit_tests/headers/headers_test_11.cu new file mode 100644 index 0000000000..7c59dbe4a7 --- /dev/null +++ b/tests/hipify-clang/unit_tests/headers/headers_test_11.cu @@ -0,0 +1,14 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args + +// Checks that HIP header file is included after include guard controlling macro, +// which goes before #pragma once. +// CHECK: #ifndef HEADERS_TEST_10_H +// CHECK-NEXT: #include +#ifndef HEADERS_TEST_10_H +// CHECK: #pragma once +#pragma once +// CHECK-NOT: #include +#define HEADERS_TEST_10_H +#include +static int counter = 0; +#endif // HEADERS_TEST_10_H diff --git a/tests/src/dynamicLoading/complex_loading_behavior.cpp b/tests/src/dynamicLoading/complex_loading_behavior.cpp index 6ee21a9cad..b8ed0f16e0 100644 --- a/tests/src/dynamicLoading/complex_loading_behavior.cpp +++ b/tests/src/dynamicLoading/complex_loading_behavior.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD_CMD: libfoo_amd %hc %S/%s -o libfoo.so -fPIC -lpthread -shared -DTEST_SHARED_LIBRARY EXCLUDE_HIP_PLATFORM nvcc + * BUILD_CMD: libfoo_amd %hc %S/%s -o libfoo.so -Xcompiler -fPIC -lpthread -shared -DTEST_SHARED_LIBRARY EXCLUDE_HIP_PLATFORM nvcc * BUILD_CMD: libfoo_nvidia %hc %S/%s -o libfoo.so -Xcompiler -fPIC -lpthread -shared -DTEST_SHARED_LIBRARY EXCLUDE_HIP_PLATFORM hcc * BUILD_CMD: %t %hc %S/%s -o %T/%t -ldl * TEST: %t diff --git a/tests/src/g++/hipMalloc.cpp b/tests/src/g++/hipMalloc.cpp index 9c1f695a85..3aab48aad5 100644 --- a/tests/src/g++/hipMalloc.cpp +++ b/tests/src/g++/hipMalloc.cpp @@ -1,7 +1,34 @@ +/* + * Copyright (c) 2019-2020 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 WARRANNTY OF ANY KIND, EXPRESS OR + * IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + * FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER + * LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + * OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + * THE SOFTWARE. + * */ + +/* HIT_START + * BUILD_CMD: hipMalloc %cxx -D__HIP_PLATFORM_HCC__ -I%hip-path/include -I/opt/rocm/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/%t -std=c++11 EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t EXCLUDE_HIP_PLATFORM nvcc + * HIT_END + */ + + #include #include int main() { int* Ad; hipMalloc((void**)&Ad, 1024); + std::cout<<"PASSED!"< #include @@ -5,4 +31,5 @@ int main() { int *Ad; hipMalloc((void**)&Ad, 1024); + printf("PASSED!\n"); } diff --git a/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp b/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp index 27947339cb..c69eb93a10 100644 --- a/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp +++ b/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp @@ -132,6 +132,10 @@ int main(int argc, char* argv[]) { CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture3DWidth, props.maxTexture3D[0])); CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture3DHeight, props.maxTexture3D[1])); CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture3DDepth, props.maxTexture3D[2])); + + CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeCooperativeLaunch, props.cooperativeLaunch)); + CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeCooperativeMultiDeviceLaunch, props.cooperativeMultiDeviceLaunch)); + #ifndef __HIP_PLATFORM_NVCC__ CHECK(test_hipDeviceGetHdpAddress(deviceId, hipDeviceAttributeHdpMemFlushCntl, props.hdpMemFlushCntl)); CHECK(test_hipDeviceGetHdpAddress(deviceId, hipDeviceAttributeHdpRegFlushCntl, props.hdpRegFlushCntl));