AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge 'master' into 'amd-master'

Browse Source 
Change-Id: Ifbc64e6a162790ce7e2e949904e9d53ef1f6540d
This commit is contained in:
Jenkins
2019-08-06 05:10:19 -04:00
parent 9fe0bd1a56 de1dd33333
commit a05c22cad3
29 changed files with 435 additions and 274 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hip-config-clang.cmake.in
+1 -14
View File
@@ -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@" )
hip-config-hcc.cmake.in
+1 -14
View File
@@ -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@" )
hip_prof_gen.py
+5 -2
View File
@@ -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' +
hipify-clang/README.md
+68 -66
View File
@@ -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
```
### <a name="windows"></a >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")
hipify-clang/src/HipifyAction.cpp
+41 -6
View File
@@ -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<clang::ASTConsumer> 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 <hip/hip_runtime.h>\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);
}
};
}
hipify-clang/src/HipifyAction.h
+4
View File
@@ -40,6 +40,7 @@ class HipifyAction : public clang::ASTFrontendAction,
public clang::ast_matchers::MatchFinder::MatchCallback {
private:
ct::Replacements* replacements;
std::map<std::string, clang::SourceLocation> Ifndefs;
std::unique_ptr<clang::ast_matchers::MatchFinder> 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.
include/hip/hcc_detail/device_functions.h
+2 -10
View File
@@ -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
include/hip/hcc_detail/hip_runtime_api.h
+26 -8
View File
@@ -45,7 +45,7 @@ THE SOFTWARE.
#include <hip/hcc_detail/hip_texture_types.h>
#include <hip/hcc_detail/hip_surface_types.h>
#if !__HIP_VDI__
#if !__HIP_VDI__ && defined(__cplusplus)
#include <hsa/hsa.h>
#include <hip/hcc_detail/program_state.hpp>
#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 <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTexture(size_t* offset, struct texture<T, dim, readMode>& 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<T, dim, readMode>& tex,
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTexture(size_t* offset, struct texture<T, dim, readMode>& 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<T, dim, readMode>& 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 <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex, hipArray_const_t array) {
return ihipBindTextureToArrayImpl(dim, readMode, array, tex.channelDesc, &tex);
return ihipBindTextureToArrayImpl(nullptr, dim, readMode, array, tex.channelDesc, &tex);
}
template <class T, int dim, enum hipTextureReadMode readMode>
hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& 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 <class T, int dim, enum hipTextureReadMode readMode>
inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode> *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
include/hip/hcc_detail/hip_vector_types.h
+21
View File
@@ -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);
include/hip/hip_runtime_api.h
+2
View File
@@ -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;
src/functional_grid_launch.inl
+2 -2
View File
@@ -47,8 +47,8 @@ namespace hip_impl
return *static_cast<hsa_agent_t*>(
stream->locked_getAv()->get_hsa_agent());
}
else if (
ihipGetTlsDefaultCtx() && ihipGetTlsDefaultCtx()->getDevice()) {
GET_TLS();
if (ihipGetTlsDefaultCtx() && ihipGetTlsDefaultCtx()->getDevice()) {
return ihipGetDevice(
ihipGetTlsDefaultCtx()->getDevice()->_deviceId)->_hsaAgent;
}
src/hip_clang.cpp
+1
View File
@@ -177,6 +177,7 @@ hipError_t hipConfigureCall(
size_t sharedMem,
hipStream_t stream)
{
GET_TLS();
auto ctx = ihipGetTlsDefaultCtx();
LockedAccessor_CtxCrit_t crit(ctx->criticalData());
src/hip_context.cpp
+18 -21
View File
@@ -29,13 +29,10 @@ THE SOFTWARE.
#include "hip_hcc_internal.h"
#include "trace_helper.h"
// Stack of contexts
thread_local std::stack<ihipCtx_t*> 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) {
src/hip_device.cpp
+9 -3
View File
@@ -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;
src/hip_error.cpp
+3 -3
View File
@@ -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) {
src/hip_hcc.cpp
+23 -40
View File
@@ -139,24 +139,12 @@ std::atomic<int> g_lastShortTid(1);
std::vector<ProfTrigger> g_dbStartTriggers;
std::vector<ProfTrigger> 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<<hip-api pid:%d tid:%s @%lu%s\n", API_COLOR, tls_tidInfo.pid(), fullStr->c_str(), apiStartTick,
fprintf(stderr, "%s<<hip-api pid:%d tid:%s @%lu%s\n", API_COLOR, tls->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<const std::runtime_error*>(&ex)) throw *rte;
if (auto lge = dynamic_cast<const std::logic_error*>(&ex)) throw *lge;
throw ex;
#else
std::cerr << ex.what() << std::endl;
src/hip_hcc_internal.h
+64 -32
View File
@@ -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<ihipCtx_t*> ctxStack;
bool getPrimaryCtx;
};
TlsData* tls_get_ptr();
#define GET_TLS() TlsData *tls = tls_get_ptr()
extern std::vector<ProfTrigger> g_dbStartTriggers;
extern std::vector<ProfTrigger> 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
src/hip_memory.cpp
+6 -6
View File
@@ -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;
src/hip_module.cpp
+25 -23
View File
@@ -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<const amd_kernel_code_v3_t*>(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<uintptr_t>(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<char> tmp{istreambuf_iterator<char>{file}, istreambuf_iterator<char>{}};
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));
}
src/hip_peer.cpp
+6 -6
View File
@@ -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));
}
src/hip_stream.cpp
+5 -5
View File
@@ -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));
}
src/hip_texture.cpp
+13 -11
View File
@@ -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);
}
src/macro_based_grid_launch.inl
+2 -1
View File
@@ -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());
}
}
}
tests/hipify-clang/unit_tests/headers/headers_test_10.cu
+14
View File
@@ -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 <hip/hip_runtime.h>
#pragma once
#ifndef HEADERS_TEST_10_H
// CHECK: #ifndef HEADERS_TEST_10_H
// CHECK-NOT: #include <hip/hip_runtime.h>
#define HEADERS_TEST_10_H
#include <stdio.h>
static int counter = 0;
#endif // HEADERS_TEST_10_H
tests/hipify-clang/unit_tests/headers/headers_test_11.cu
+14
View File
@@ -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 <hip/hip_runtime.h>
#ifndef HEADERS_TEST_10_H
// CHECK: #pragma once
#pragma once
// CHECK-NOT: #include <hip/hip_runtime.h>
#define HEADERS_TEST_10_H
#include <stdio.h>
static int counter = 0;
#endif // HEADERS_TEST_10_H
tests/src/dynamicLoading/complex_loading_behavior.cpp
+1 -1
View File
@@ -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
tests/src/g++/hipMalloc.cpp
+27
View File
@@ -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 <hip/hip_runtime_api.h>
#include <iostream>
int main() {
int* Ad;
hipMalloc((void**)&Ad, 1024);
std::cout<<"PASSED!"<<std::endl;
}
tests/src/gcc/hipMalloc.c
+27
View File
@@ -1,3 +1,29 @@
/*
* 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 %cc -D__HIP_PLATFORM_HCC__ -I%hip-path/include %S/%s -Wl,--rpath=%hip-path/lib %hip-path/lib/libhip_hcc.so -o %T/%t EXCLUDE_HIP_PLATFORM nvcc
* TEST: %t EXCLUDE_HIP_PLATFORM nvcc
* HIT_END
*/
#include<hip/hip_runtime_api.h>
#include<stdio.h>
@@ -5,4 +31,5 @@ int main()
{
int *Ad;
hipMalloc((void**)&Ad, 1024);
printf("PASSED!\n");
}
tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp
+4
View File
@@ -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));