diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt index a26e848985..37d3534cd8 100644 --- a/projects/hip/CMakeLists.txt +++ b/projects/hip/CMakeLists.txt @@ -2,18 +2,45 @@ cmake_minimum_required(VERSION 2.8.3) project(hip) ############################# -# Configure variables +# Setup config generation ############################# -# Determine HIP_VERSION +string(TIMESTAMP _timestamp UTC) +set(_versionInfo "# Auto-generated by cmake\n") +set(_buildInfo "# Auto-generated by cmake on ${_timestamp} UTC\n") +macro(add_to_config _configfile _variable) + set(${_configfile} "${${_configfile}}${_variable}=${${_variable}}\n") +endmacro() + +############################# +# Setup version information +############################# +# Determine HIP_BASE_VERSION execute_process(COMMAND ${CMAKE_CURRENT_SOURCE_DIR}/bin/hipconfig --version - OUTPUT_VARIABLE HIP_VERSION + OUTPUT_VARIABLE HIP_BASE_VERSION OUTPUT_STRIP_TRAILING_WHITESPACE) -string(REPLACE "." ";" VERSION_LIST ${HIP_VERSION}) +string(REPLACE "." ";" VERSION_LIST ${HIP_BASE_VERSION}) list(GET VERSION_LIST 0 HIP_VERSION_MAJOR) list(GET VERSION_LIST 1 HIP_VERSION_MINOR) -list(GET VERSION_LIST 2 HIP_VERSION_PATCH) +# get date information based on UTC +# use the last two digits of year + week number + day in the week as HIP_VERSION_PATCH +# use the commit date, instead of build date +# add xargs to remove strange trailing newline character +execute_process(COMMAND git show -s --format=@%ct + COMMAND xargs + COMMAND date -f - --utc +%y%U%w + WORKING_DIRECTORY ${PROJECT_SOURCE_DIR} + OUTPUT_VARIABLE HIP_VERSION_PATCH + OUTPUT_STRIP_TRAILING_WHITESPACE) +set(HIP_VERSION $HIP_VERSION_MAJOR.$HIP_VERSION_MINOR.$HIP_VERSION_PATCH) +add_to_config(_versionInfo HIP_VERSION_MAJOR) +add_to_config(_versionInfo HIP_VERSION_MINOR) +add_to_config(_versionInfo HIP_VERSION_PATCH) + +############################# +# Configure variables +############################# # Determine HIP_PLATFORM if(NOT DEFINED HIP_PLATFORM) if(NOT DEFINED ENV{HIP_PLATFORM}) @@ -36,6 +63,9 @@ if(HIP_PLATFORM STREQUAL "hcc") set(HCC_HOME $ENV{HCC_HOME} CACHE PATH "Path to which HCC has been installed") endif() endif() + if(DEFINED ENV{HIP_DEVELOPER}) + add_to_config(_buildInfo HCC_HOME) + endif() if(IS_ABSOLUTE ${HCC_HOME} AND EXISTS ${HCC_HOME} AND IS_DIRECTORY ${HCC_HOME}) execute_process(COMMAND ${HCC_HOME}/bin/hcc --version COMMAND cut -d\ -f9 @@ -45,6 +75,7 @@ if(HIP_PLATFORM STREQUAL "hcc") else() message(FATAL_ERROR "Don't know where to find HCC. Please specify abolute path using -DHCC_HOME") endif() + add_to_config(_buildInfo HCC_VERSION) # Determine HSA_PATH if(NOT DEFINED HSA_PATH) @@ -83,11 +114,14 @@ else() endif() # Set if we need to build shared or static library -if(NOT DEFINED ENV{HIP_USE_SHARED_LIBRARY}) - set(HIP_USE_SHARED_LIBRARY 0) -else() - set(HIP_USE_SHARED_LIBRARY $ENV{HIP_USE_SHARED_LIBRARY}) +if(NOT DEFINED HIP_LIB_TYPE) + if(NOT DEFINED ENV{HIP_LIB_TYPE}) + set(HIP_LIB_TYPE 0) + else() + set(HIP_LIB_TYPE $ENV{HIP_LIB_TYPE}) + endif() endif() +add_to_config(_buildInfo HIP_LIB_TYPE) # Check if we need to build clang hipify if(NOT DEFINED BUILD_CLANG_HIPIFY) @@ -106,6 +140,7 @@ if(NOT DEFINED COMPILE_HIP_ATP_MARKER) set(COMPILE_HIP_ATP_MARKER $ENV{COMPILE_HIP_ATP_MARKER}) endif() endif() +add_to_config(_buildInfo COMPILE_HIP_ATP_MARKER) ############################# # Build steps @@ -133,6 +168,7 @@ if(HIP_PLATFORM STREQUAL "hcc") set(SOURCE_FILES src/device_util.cpp src/hip_hcc.cpp + src/hip_context.cpp src/hip_device.cpp src/hip_error.cpp src/hip_event.cpp @@ -141,20 +177,24 @@ if(HIP_PLATFORM STREQUAL "hcc") src/hip_peer.cpp src/hip_stream.cpp src/hip_fp16.cpp - src/staging_buffer.cpp) + src/unpinned_copy_engine.cpp + src/hip_module.cpp) - if(${HIP_USE_SHARED_LIBRARY} EQUAL 1) - add_library(hip_hcc SHARED ${SOURCE_FILES}) - else() - #add_library(hip_hcc STATIC ${SOURCE_FILES}) + if(${HIP_LIB_TYPE} EQUAL 0) add_library(hip_hcc OBJECT ${SOURCE_FILES}) + elseif(${HIP_LIB_TYPE} EQUAL 1) + add_library(hip_hcc STATIC ${SOURCE_FILES}) + else() + add_library(hip_hcc SHARED ${SOURCE_FILES}) endif() - # Generate .hipconfig - string(TIMESTAMP _timestamp) - file(WRITE "${PROJECT_BINARY_DIR}/.hipconfig" "# Auto-generated by cmake on ${_timestamp} local time\nHCC_HOME=${HCC_HOME}\nHCC_VERSION=${HCC_VERSION}\n") + # Generate .buildInfo + file(WRITE "${PROJECT_BINARY_DIR}/.buildInfo" ${_buildInfo}) endif() +# Generate .version +file(WRITE "${PROJECT_BINARY_DIR}/.version" ${_versionInfo}) + # Build doxygen documentation add_custom_target(doc COMMAND HIP_PATH=${CMAKE_CURRENT_SOURCE_DIR} doxygen ${CMAKE_CURRENT_SOURCE_DIR}/docs/doxygen-input/doxy.cfg WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR}/docs) @@ -164,17 +204,19 @@ add_custom_target(doc COMMAND HIP_PATH=${CMAKE_CURRENT_SOURCE_DIR} doxygen ${CMA ############################# # Install hip_hcc if platform is hcc if(HIP_PLATFORM STREQUAL "hcc") - if(${HIP_USE_SHARED_LIBRARY} EQUAL 1) - install(TARGETS hip_hcc DESTINATION lib) - else() - #install(TARGETS hip_hcc DESTINATION lib) + if(${HIP_LIB_TYPE} EQUAL 0) install(DIRECTORY ${PROJECT_BINARY_DIR}/CMakeFiles/hip_hcc.dir/src/ DESTINATION lib) + else() + install(TARGETS hip_hcc DESTINATION lib) endif() - # Install .hipconfig - install(FILES ${PROJECT_BINARY_DIR}/.hipconfig DESTINATION lib) + # Install .buildInfo + install(FILES ${PROJECT_BINARY_DIR}/.buildInfo DESTINATION lib) endif() +# Install .version +install(FILES ${PROJECT_BINARY_DIR}/.version DESTINATION bin) + # Install src, bin, include if necessary execute_process(COMMAND test ${CMAKE_INSTALL_PREFIX} -ef ${CMAKE_CURRENT_SOURCE_DIR} RESULT_VARIABLE INSTALL_SOURCE) diff --git a/projects/hip/CONTRIBUTING.md b/projects/hip/CONTRIBUTING.md index f6e578efd8..cd785fe336 100644 --- a/projects/hip/CONTRIBUTING.md +++ b/projects/hip/CONTRIBUTING.md @@ -90,6 +90,33 @@ The HIP interface is designed to be very familiar for CUDA programmers. Differences or limitations of HIP APIs as compared to CUDA APIs should be clearly documented and described. +## Coding Guidelines (in brief) +- Code Indentation: + - Tabs should be expanded to spaces. + - Use 4 spaces indendation. +- Capitaliziation and Naming + - Prefer camelCase for HIP interfaces and internal symbols. Note HCC uses _ for separator. + This guideline is not yet consistently followed in HIP code - eventual compliance is aspirational. + - Member variables should begin with a leading "_". This allows them to be easily distinguished from other variables or functions. + + +- {} placement + - For functions, the opening { should be placed on a new line. + - For if/else blocks, the opening { is placed on same line as the if/else. Use a space to separate {/" from if/else. Example +''' + if (foo) { + doFoo() + } else { + doFooElse(); + } +''' + - Single-line if statement should still use {/} pair (even though C++ does not require). +- Miscellaneous + - All references in function parameter lists should be const. + - "ihip" = internal hip structures. These should not be exposed through the HIP API. + - Keyword TODO refers to a note that should be addressed in long-term. Could be style issue, software architecture, or known bugs. + - FIXME refers to a short-term bug that needs to be addressed. + #### Presubmit Testing: Before checking in or submitting a pull request, run all Rodinia tests and ensure pass results match starting point: diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 1141e1c08d..c51e351382 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -1,5 +1,8 @@ #!/usr/bin/perl -w + +# Need perl > 5.10 to use logic-defined or +use 5.006; use v5.10.1; use File::Basename; # HIP compiler driver @@ -24,39 +27,40 @@ print "No Arguments passed, exiting ...\n"; exit(-1); } -$verbose = $ENV{'HIPCC_VERBOSE'}; -$verbose = 0 unless defined $verbose; +#--- +# Function to parse config file +sub parse_config_file { + my ($file, $config) = @_; + if (open (CONFIG, "$file")) { + while () { + my $config_line=$_; + chop ($config_line); + $config_line =~ s/^\s*//; + $config_line =~ s/\s*$//; + if (($config_line !~ /^#/) && ($config_line ne "")) { + my ($name, $value) = split (/=/, $config_line); + $$config{$name} = $value; + } + } + close(CONFIG); + } +} + +$verbose = $ENV{'HIPCC_VERBOSE'} // 0; # Verbose: 0x1=commands, 0x2=paths, 0x4=hippc args -$HIP_PATH=$ENV{'HIP_PATH'}; -$HIP_PATH=dirname (dirname $0) unless defined $HIP_PATH; # use parent directory of hipcc +$HIP_PATH=$ENV{'HIP_PATH'} // dirname (dirname $0); # use parent directory of hipcc #--- -# Read .hipconfig +# Read .buildInfo my %hipConfig = (); -$hipConfig{'VALID'}=0; -if (open (CONFIG, "$HIP_PATH/lib/.hipconfig")) { - while () { - my $config_line=$_; - chop ($config_line); - $config_line =~ s/^\s*//; - $config_line =~ s/\s*$//; - if (($config_line !~ /^#/) && ($config_line ne "")) { - my ($name, $value) = split (/=/, $config_line); - $hipConfig{$name} = $value; - $hipConfig{'VALID'}=1; - } - } - close(CONFIG); -} +parse_config_file("$HIP_PATH/lib/.buildInfo", \%hipConfig); #--- #HIP_PLATFORM controls whether to use NVCC or HCC for compilation: -$HIP_PLATFORM= `$HIP_PATH/bin/hipconfig --platform`; +$HIP_PLATFORM= `$HIP_PATH/bin/hipconfig --platform` // "hcc"; $HIP_VERSION= `$HIP_PATH/bin/hipconfig --version`; -$HIP_PLATFORM="hcc" unless defined $HIP_PLATFORM; - if ($verbose & 0x2) { print ("HIP_PATH=$HIP_PATH\n"); print ("HIP_PLATFORM=$HIP_PLATFORM\n"); @@ -66,21 +70,18 @@ if ($verbose & 0x2) { $setStdLib = 0; # TODO - set to 0 if ($HIP_PLATFORM eq "hcc") { - $HSA_PATH=$ENV{'HSA_PATH'}; - $HSA_PATH="/opt/rocm/hsa" unless defined $HSA_PATH; + $HSA_PATH=$ENV{'HSA_PATH'} // "/opt/rocm/hsa"; + + $HCC_HOME=$ENV{'HCC_HOME'} // $hipConfig{'HCC_HOME'} // "/opt/rocm/hcc"; - $HCC_HOME=$ENV{'HCC_HOME'}; - $HCC_HOME="/opt/rocm/hcc" unless defined $HCC_HOME; $HCC_VERSION=`${HCC_HOME}/bin/hcc --version | cut -d" " -f9 | tr -d "\n"`; - $ROCM_PATH=$ENV{'ROCM_PATH'}; - $ROCM_PATH="/opt/rocm" unless defined $ROCM_PATH; + $ROCM_PATH=$ENV{'ROCM_PATH'} // "/opt/rocm"; $HIP_ATP_MARKER=$ENV{'HIP_ATP_MARKER'}; $marker_path = "$ROCM_PATH/profiler/CXLActivityLogger"; - $ROCM_TARGET=$ENV{'ROCM_TARGET'}; - $ROCM_TARGET="fiji" unless defined $ROCM_TARGET; + $ROCM_TARGET=$ENV{'ROCM_TARGET'} // "fiji"; # HCC* may be used to compile src/hip_hcc.o (and also feed the HIPCXXFLAGS below) $HCC = "$HCC_HOME/bin/hcc"; @@ -111,6 +112,9 @@ if ($HIP_PLATFORM eq "hcc") { if ($ROCM_TARGET eq "fiji") { $HIPLDFLAGS .= " -amdgpu-target=AMD:AMDGPU:8:0:3"; } + if ($ROCM_TARGET eq "carrizo") { + $HIPLDFLAGS .= " -amdgpu-target=AMD:AMDGPU:8:0:1"; + } if ($ROCM_TARGET eq "hawaii") { $HIPLDFLAGS .= " -amdgpu-target=AMD:AMDGPU:7:0:1"; } @@ -142,8 +146,7 @@ if ($HIP_PLATFORM eq "hcc") { if ($verbose & 0x2) { print ("CUDA_PATH=$CUDA_PATH\n"); } - $CUDA_PATH=$ENV{'CUDA_PATH'}; - $CUDA_PATH='/usr/local/cuda' unless defined $CUDA_PATH; + $CUDA_PATH=$ENV{'CUDA_PATH'} // '/usr/local/cuda'; $HIPCC="$CUDA_PATH/bin/nvcc"; $HIPCXXFLAGS .= " -I$CUDA_PATH/include"; @@ -249,14 +252,14 @@ if ($setStdLib eq 0 and $HIP_PLATFORM eq 'hcc') } if ($needHipHcc) { - $HIP_USE_SHARED_LIBRARY = $ENV{'HIP_USE_SHARED_LIBRARY'}; - $HIP_USE_SHARED_LIBRARY = 0 unless defined $HIP_USE_SHARED_LIBRARY; + $HIP_LIB_TYPE = $hipConfig{'HIP_LIB_TYPE'} // 0; - #$HIPLDFLAGS .= " -L/opt/rocm/hip/lib -lhip_hcc" ; - if ($HIP_USE_SHARED_LIBRARY) { - $HIPLDFLAGS .= " -L$HIP_PATH/lib -Wl,--rpath=$HIP_PATH/lib -lhip_hcc"; + if ($HIP_LIB_TYPE eq 0) { + $HIPLDFLAGS .= " $HIP_PATH/lib/device_util.cpp.o $HIP_PATH/lib/hip_device.cpp.o $HIP_PATH/lib/hip_error.cpp.o $HIP_PATH/lib/hip_event.cpp.o $HIP_PATH/lib/hip_hcc.cpp.o $HIP_PATH/lib/hip_memory.cpp.o $HIP_PATH/lib/hip_peer.cpp.o $HIP_PATH/lib/hip_stream.cpp.o $HIP_PATH/lib/unpinned_copy_engine.cpp.o $HIP_PATH/lib/hip_ldg.cpp.o $HIP_PATH/lib/hip_fp16.cpp.o $HIP_PATH/lib/hip_context.cpp.o $HIP_PATH/lib/hip_module.cpp.o"; + } elsif ($HIP_LIB_TYPE eq 1) { + $HIPLDFLAGS .= " -L$HIP_PATH/lib -lhip_hcc" ; } else { - $HIPLDFLAGS .= " $HIP_PATH/lib/device_util.cpp.o $HIP_PATH/lib/hip_device.cpp.o $HIP_PATH/lib/hip_error.cpp.o $HIP_PATH/lib/hip_event.cpp.o $HIP_PATH/lib/hip_hcc.cpp.o $HIP_PATH/lib/hip_memory.cpp.o $HIP_PATH/lib/hip_peer.cpp.o $HIP_PATH/lib/hip_stream.cpp.o $HIP_PATH/lib/staging_buffer.cpp.o $HIP_PATH/lib/hip_ldg.cpp.o $HIP_PATH/lib/hip_fp16.cpp.o"; + $HIPLDFLAGS .= " -L$HIP_PATH/lib -Wl,--rpath=$HIP_PATH/lib -lhip_hcc"; } } @@ -287,7 +290,7 @@ if ($printHipVersion) { print $HIP_VERSION, "\n"; } if ($runCmd) { - if ($hipConfig{'VALID'} and $HIP_PLATFORM eq "hcc" and $HCC_VERSION ne $hipConfig{'HCC_VERSION'}) { + if ($HIP_PLATFORM eq "hcc" and exists($hipConfig{'HCC_VERSION'}) and $HCC_VERSION ne $hipConfig{'HCC_VERSION'}) { print ("HIP was built using $hipConfig{'HCC_VERSION'}, but you are using $HCC_VERSION. Please rebuild HIP.\n") && die (); } system ("$CMD") and die (); diff --git a/projects/hip/bin/hipconfig b/projects/hip/bin/hipconfig index 4f65d3230c..4fc37944e8 100755 --- a/projects/hip/bin/hipconfig +++ b/projects/hip/bin/hipconfig @@ -1,9 +1,10 @@ #!/usr/bin/perl -w -$HIP_VERSION_MAJOR = "0"; -$HIP_VERSION_MINOR = "92"; -$HIP_VERSION_PATCH = "0"; +$HIP_BASE_VERSION_MAJOR = "0"; +$HIP_BASE_VERSION_MINOR = "92"; +# Need perl > 5.10 to use logic-defined or +use 5.006; use v5.10.1; use Getopt::Long; use Cwd; @@ -39,23 +40,35 @@ if ($p_help) { exit(); } -$HIP_VERSION="$HIP_VERSION_MAJOR.$HIP_VERSION_MINOR.$HIP_VERSION_PATCH"; +#--- +# Function to parse config file +sub parse_config_file { + my ($file, $config) = @_; + if (open (CONFIG, "$file")) { + while () { + my $config_line=$_; + chop ($config_line); + $config_line =~ s/^\s*//; + $config_line =~ s/\s*$//; + if (($config_line !~ /^#/) && ($config_line ne "")) { + my ($name, $value) = split (/=/, $config_line); + $$config{$name} = $value; + } + } + close(CONFIG); + } +} -$CUDA_PATH=$ENV{'CUDA_PATH'}; -$CUDA_PATH='/usr/local/cuda' unless defined $CUDA_PATH; - -$HCC_HOME=$ENV{'HCC_HOME'}; -$HCC_HOME='/opt/rocm/hcc' unless defined $HCC_HOME; - -$HSA_PATH=$ENV{'HSA_PATH'}; -$HSA_PATH='/opt/rocm/hsa' unless defined $HSA_PATH; +$CUDA_PATH=$ENV{'CUDA_PATH'} // '/usr/local/cuda'; +$HCC_HOME=$ENV{'HCC_HOME'} // '/opt/rocm/hcc'; +$HSA_PATH=$ENV{'HSA_PATH'} // '/opt/rocm/hsa'; #--- #HIP_PLATFORM controls whether to use NVCC or HCC for compilation: $HIP_PLATFORM=$ENV{'HIP_PLATFORM'}; if (not defined $HIP_PLATFORM) { $NAMDGPUNODES=`cat /sys/class/kfd/kfd/topology/nodes/*/properties 2>/dev/null | grep -c 'simd_count [1-9]'`; - + if ($NAMDGPUNODES > 0) { $HIP_PLATFORM = "hcc" } else { @@ -63,9 +76,7 @@ if (not defined $HIP_PLATFORM) { } } -$HIP_PATH=$ENV{'HIP_PATH'}; -$HIP_PATH=Cwd::realpath (dirname (dirname $0)) unless defined $HIP_PATH; # use parent directory of this tool - +$HIP_PATH=$ENV{'HIP_PATH'} // Cwd::realpath (dirname (dirname $0)); # use parent directory of this tool if ($HIP_PLATFORM eq "hcc") { $CPP_CONFIG= " -D__HIP_PLATFORM_HCC__= -I$HIP_PATH/include -I$HCC_HOME/include"; @@ -74,11 +85,20 @@ if ($HIP_PLATFORM eq "nvcc") { $CPP_CONFIG = " -D__HIP_PLATFORM_NVCC__= -I$HIP_PATH/include -I$CUDA_PATH/include"; }; +#--- +# Read .version +my %hipVersion = (); +parse_config_file("$HIP_PATH/bin/.version", \%hipVersion); +$HIP_VERSION_MAJOR = $hipVersion{'HIP_VERSION_MAJOR'} // $HIP_BASE_VERSION_MAJOR; +$HIP_VERSION_MINOR = $hipVersion{'HIP_VERSION_MINOR'} // $HIP_BASE_VERSION_MINOR; +$HIP_VERSION_PATCH = $hipVersion{'HIP_VERSION_PATCH'} // "0"; +$HIP_VERSION="$HIP_VERSION_MAJOR.$HIP_VERSION_MINOR.$HIP_VERSION_PATCH"; + if ($p_path) { print "$HIP_PATH"; $printed = 1; } - + if ($p_cpp_config) { print $CPP_CONFIG; @@ -102,16 +122,16 @@ if (!$printed or $p_full) { print "HIP_PATH : ", $HIP_PATH, "\n"; print "HIP_PLATFORM : ", $HIP_PLATFORM, "\n"; print "CPP_CONFIG : ", $CPP_CONFIG, "\n"; - if ($HIP_PLATFORM eq "hcc") + if ($HIP_PLATFORM eq "hcc") { print "\n" ; print "== hcc\n"; print ("HSA_PATH : $HSA_PATH\n"); print ("HCC_HOME : $HCC_HOME\n"); system("$HCC_HOME/bin/hcc --version"); - print ("HCC-cxxflags : "); + print ("HCC-cxxflags : "); system("$HCC_HOME/bin/hcc-config --cxxflags"); - print ("HCC-ldflags : "); + print ("HCC-ldflags : "); system("$HCC_HOME/bin/hcc-config --ldflags"); printf("\n"); } diff --git a/projects/hip/clang-hipify/src/Cuda2Hip.cpp b/projects/hip/clang-hipify/src/Cuda2Hip.cpp index 02edeef95b..45960ac8ca 100644 --- a/projects/hip/clang-hipify/src/Cuda2Hip.cpp +++ b/projects/hip/clang-hipify/src/Cuda2Hip.cpp @@ -69,15 +69,17 @@ enum ConvTypes { CONV_TEX, CONV_OTHER, CONV_INCLUDE, + CONV_INCLUDE_CUDA_MAIN_H, CONV_LITERAL, CONV_BLAS, CONV_LAST }; const char *counterNames[ConvTypes::CONV_LAST] = { - "dev", "mem", "kern", "coord_func", "math_func", - "special_func", "stream", "event", "err", "def", - "tex", "other", "include", "literal", "blas"}; + "dev", "mem", "kern", "coord_func", "math_func", + "special_func", "stream", "event", "err", "def", + "tex", "other", "include", "include_cuda_main_header", + "literal", "blas"}; namespace { @@ -91,7 +93,8 @@ struct cuda2hipMap { cuda2hipRename["__CUDACC__"] = {"__HIPCC__", CONV_DEF}; // CUDA includes - cuda2hipRename["cuda_runtime.h"] = {"hip_runtime.h", CONV_INCLUDE}; + cuda2hipRename["cuda.h"] = {"hip_runtime.h", CONV_INCLUDE_CUDA_MAIN_H}; + cuda2hipRename["cuda_runtime.h"] = {"hip_runtime.h", CONV_INCLUDE_CUDA_MAIN_H}; cuda2hipRename["cuda_runtime_api.h"] = {"hip_runtime_api.h", CONV_INCLUDE}; // HIP includes @@ -1040,7 +1043,10 @@ static void processString(StringRef s, const cuda2hipMap &map, } } -struct HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks { +class Cuda2HipCallback; + +class HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks { +public: HipifyPPCallbacks(Replacements *R) : SeenEnd(false), _sm(nullptr), _pp(nullptr), Replace(R) {} @@ -1055,6 +1061,8 @@ struct HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks { return true; } + virtual void handleEndSource() override; + virtual void InclusionDirective(SourceLocation hash_loc, const Token &include_token, StringRef file_name, bool is_angled, @@ -1151,11 +1159,29 @@ struct HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks { Replacement Rep(*_sm, sl, name.size(), repName); Replace->insert(Rep); } - } - if (tok.is(tok::string_literal)) { - StringRef s(tok.getLiteralData(), tok.getLength()); - processString(unquoteStr(s), N, Replace, *_sm, tok.getLocation(), - countReps); + } else if (tok.isLiteral()) { + SourceLocation sl = tok.getLocation(); + if (_sm->isMacroBodyExpansion(sl)) { + LangOptions DefaultLangOptions; + SourceLocation sl_macro = _sm->getExpansionLoc(sl); + SourceLocation sl_end = Lexer::getLocForEndOfToken(sl_macro, 0, *_sm, DefaultLangOptions); + size_t length = _sm->getCharacterData(sl_end) - _sm->getCharacterData(sl_macro); + StringRef name = StringRef(_sm->getCharacterData(sl_macro), length); + const auto found = N.cuda2hipRename.find(name); + if (found != N.cuda2hipRename.end()) { + sl = sl_macro; + countReps[found->second.countType]++; + StringRef repName = found->second.hipName; + Replacement Rep(*_sm, sl, length, repName); + Replace->insert(Rep); + } + } else { + if (tok.is(tok::string_literal)) { + StringRef s(tok.getLiteralData(), tok.getLength()); + processString(unquoteStr(s), N, Replace, *_sm, tok.getLocation(), + countReps); + } + } } } } @@ -1168,21 +1194,23 @@ struct HipifyPPCallbacks : public PPCallbacks, public SourceFileCallbacks { bool SeenEnd; void setSourceManager(SourceManager *sm) { _sm = sm; } void setPreprocessor(Preprocessor *pp) { _pp = pp; } - + void setMatch(Cuda2HipCallback *match) { Match = match; } int64_t countReps[ConvTypes::CONV_LAST] = {0}; private: SourceManager *_sm; Preprocessor *_pp; - + Cuda2HipCallback *Match; Replacements *Replace; struct cuda2hipMap N; }; class Cuda2HipCallback : public MatchFinder::MatchCallback { public: - Cuda2HipCallback(Replacements *Replace, ast_matchers::MatchFinder *parent) - : Replace(Replace), owner(parent) {} + Cuda2HipCallback(Replacements *Replace, ast_matchers::MatchFinder *parent, HipifyPPCallbacks *PPCallbacks) + : Replace(Replace), owner(parent), PP(PPCallbacks) { + PP->setMatch(this); + } void convertKernelDecl(const FunctionDecl *kernelDecl, const MatchFinder::MatchResult &Result) { @@ -1424,6 +1452,42 @@ public: } } + if (const VarDecl *sharedVar = + Result.Nodes.getNodeAs("cudaSharedIncompleteArrayVar")) { + // Example: extern __shared__ uint sRadix1[]; + if (sharedVar->hasExternalFormalLinkage()) { + QualType QT = sharedVar->getType(); + StringRef typeName; + if (QT->isIncompleteArrayType()) { + const ArrayType *AT = QT.getTypePtr()->getAsArrayTypeUnsafe(); + QT = AT->getElementType(); + if (QT.getTypePtr()->isBuiltinType()) { + QT = QT.getCanonicalType(); + const BuiltinType *BT = dyn_cast(QT); + if (BT) { + LangOptions LO; + LO.CUDA = true; + PrintingPolicy policy(LO); + typeName = BT->getName(policy); + } + } else { + typeName = QT.getAsString(); + } + } + if (!typeName.empty()) { + SourceLocation slStart = sharedVar->getLocStart(); + SourceLocation slEnd = sharedVar->getLocEnd(); + size_t repLength = SM->getCharacterData(slEnd) - SM->getCharacterData(slStart) + 1; + SmallString<128> tmpData; + StringRef varName = sharedVar->getNameAsString(); + StringRef repName = Twine("HIP_DYNAMIC_SHARED(" + typeName + ", " + varName + ")").toStringRef(tmpData); + Replacement Rep(*SM, slStart, repLength, repName); + Replace->insert(Rep); + countReps[CONV_MEM]++; + } + } + } + if (const VarDecl *cudaStructVarPtr = Result.Nodes.getNodeAs("cudaStructVarPtr")) { const Type *t = cudaStructVarPtr->getType().getTypePtrOrNull(); @@ -1507,6 +1571,14 @@ public: Replace->insert(Rep); } } + + if (PP->countReps[CONV_INCLUDE_CUDA_MAIN_H] == 0 && + countReps[CONV_INCLUDE_CUDA_MAIN_H] == 0 && Replace->size() > 0) { + StringRef repName = "#include \n"; + Replacement Rep(*SM, SM->getLocForStartOfFile(SM->getMainFileID()), 0, repName); + Replace->insert(Rep); + countReps[CONV_INCLUDE_CUDA_MAIN_H]++; + } } int64_t countReps[ConvTypes::CONV_LAST] = {0}; @@ -1514,9 +1586,20 @@ public: private: Replacements *Replace; ast_matchers::MatchFinder *owner; + HipifyPPCallbacks *PP; struct cuda2hipMap N; }; +void HipifyPPCallbacks::handleEndSource() { + if (Match->countReps[CONV_INCLUDE_CUDA_MAIN_H] == 0 && + countReps[CONV_INCLUDE_CUDA_MAIN_H] == 0 && Replace->size() > 0) { + StringRef repName = "#include \n"; + Replacement Rep(*_sm, _sm->getLocForStartOfFile(_sm->getMainFileID()), 0, repName); + Replace->insert(Rep); + countReps[CONV_INCLUDE_CUDA_MAIN_H]++; + } +} + } // end anonymous namespace // Set up the command line options @@ -1578,8 +1661,9 @@ int main(int argc, const char **argv) { RefactoringTool Tool(OptionsParser.getCompilations(), dst); ast_matchers::MatchFinder Finder; - Cuda2HipCallback Callback(&Tool.getReplacements(), &Finder); HipifyPPCallbacks PPCallbacks(&Tool.getReplacements()); + Cuda2HipCallback Callback(&Tool.getReplacements(), &Finder, &PPCallbacks); + Finder.addMatcher(callExpr(isExpansionInMainFile(), callee(functionDecl(matchesName("cuda.*|cublas.*")))) .bind("cudaCall"), @@ -1634,6 +1718,11 @@ int main(int argc, const char **argv) { &Callback); Finder.addMatcher(stringLiteral(isExpansionInMainFile()).bind("stringLiteral"), &Callback); + Finder.addMatcher(varDecl(isExpansionInMainFile(), allOf( + hasAttr(attr::CUDAShared), + hasType(incompleteArrayType()))) + .bind("cudaSharedIncompleteArrayVar"), + &Callback); auto action = newFrontendActionFactory(&Finder, &PPCallbacks); diff --git a/projects/hip/docs/markdown/hip_faq.md b/projects/hip/docs/markdown/hip_faq.md index 0f2a43fe26..8b50a9a8c7 100644 --- a/projects/hip/docs/markdown/hip_faq.md +++ b/projects/hip/docs/markdown/hip_faq.md @@ -208,3 +208,9 @@ HIP_TRACE_API=1 HIP_DB=0x2 ./myHipApp ``` Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors. + +### What if HIP generates error of "symbol multiply defined!" only on AMD machine? +Unlike CUDA, in HCC, for functions defined in the header files, the keyword of "__forceinline__" does not imply "static". +Thus, if failed to define "static" keyword, you might see a lot of "symbol multiply defined!" errors at compilation. +The workaround is to explicitly add the keyword of "static" before any functions that were defined as "__forceinline__". + diff --git a/projects/hip/docs/markdown/hip_porting_guide.md b/projects/hip/docs/markdown/hip_porting_guide.md index 7857e4b983..4d376f4a5f 100644 --- a/projects/hip/docs/markdown/hip_porting_guide.md +++ b/projects/hip/docs/markdown/hip_porting_guide.md @@ -264,6 +264,38 @@ Makefiles can use the following syntax to conditionally provide a default HIP_PA HIP_PATH ?= $(shell hipconfig --path) ``` +## hipLaunchKernel + +hipLaunchKernel is a variadic macro which accepts as parameters the launch configurations (grid dims, group dims, stream, dynamic shared size) followed by a variable number of kernel arguments. +This sequence is then expanded into the appropriate kernel launch syntax depending on the platform. +While this can be a convenient single-line kernel launch syntax, the macro implementation can cause issues when nested inside other macros. For example, consider the following: + +``` +// Will cause compile error: +#define MY_LAUNCH(command, doTrace) \ +{\ + if (doTrace) printf ("TRACE: %s\n", #command); \ + (command); /* The nested ( ) will cause compile error */\ +} + +MY_LAUNCH (hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCall"); +``` + +Avoid nesting macro parameters inside parenthesis - here's an alternative that will work: + +``` +#define MY_LAUNCH(command, doTrace) \ +{\ + if (doTrace) printf ("TRACE: %s\n", #command); \ + command;\ +} + +MY_LAUNCH (hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCall"); +``` + + + + ## Compiler Options hipcc is a portable compiler driver that will call nvcc or hcc (depending on the target system) and attach all required include and library options. It passes options through to the target compiler. Tools that call hipcc must ensure the compiler options are appropriate for the target compiler. The `hipconfig` script may helpful in making diff --git a/projects/hip/include/hcc_detail/hip_blas.h b/projects/hip/include/hcc_detail/hip_blas.h new file mode 100644 index 0000000000..07f41ec71b --- /dev/null +++ b/projects/hip/include/hcc_detail/hip_blas.h @@ -0,0 +1,258 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#pragma once + +#include +#include + +//HGSOS for Kalmar leave it as C++, only cublas needs C linkage. + +#ifdef __cplusplus +extern "C" { +#endif + +typedef hcblasHandle_t* hipblasHandle_t; +typedef hcComplex hipComplex ; + +static hipblasHandle_t dummyGlobal; + +/* Unsupported types + "cublasFillMode_t", + "cublasDiagType_t", + "cublasSideMode_t", + "cublasPointerMode_t", + "cublasAtomicsMode_t", + "cublasDataType_t" +*/ + +inline static hcblasOperation_t hipOperationToHCCOperation( hipblasOperation_t op) +{ + switch (op) + { + case HIPBLAS_OP_N: + return HCBLAS_OP_N; + + case HIPBLAS_OP_T: + return HCBLAS_OP_T; + + case HIPBLAS_OP_C: + return HCBLAS_OP_C; + + default: + throw "Non existent OP"; + } +} + +inline static hipblasOperation_t HCCOperationToHIPOperation( hcblasOperation_t op) +{ + switch (op) + { + case HCBLAS_OP_N : + return HIPBLAS_OP_N; + + case HCBLAS_OP_T : + return HIPBLAS_OP_T; + + case HCBLAS_OP_C : + return HIPBLAS_OP_C; + + default: + throw "Non existent OP"; + } +} + + +inline static hipblasStatus_t hipHCBLASStatusToHIPStatus(hcblasStatus_t hcStatus) +{ + switch(hcStatus) + { + case HCBLAS_STATUS_SUCCESS: + return HIPBLAS_STATUS_SUCCESS; + case HCBLAS_STATUS_NOT_INITIALIZED: + return HIPBLAS_STATUS_NOT_INITIALIZED; + case HCBLAS_STATUS_ALLOC_FAILED: + return HIPBLAS_STATUS_ALLOC_FAILED; + case HCBLAS_STATUS_INVALID_VALUE: + return HIPBLAS_STATUS_INVALID_VALUE; + case HCBLAS_STATUS_MAPPING_ERROR: + return HIPBLAS_STATUS_MAPPING_ERROR; + case HCBLAS_STATUS_EXECUTION_FAILED: + return HIPBLAS_STATUS_EXECUTION_FAILED; + case HCBLAS_STATUS_INTERNAL_ERROR: + return HIPBLAS_STATUS_INTERNAL_ERROR; + default: + throw "Unimplemented status"; + } +} + + + +inline static hipblasStatus_t hipblasCreate(hipblasHandle_t* handle) { + hipblasStatus_t retVal = hipHCBLASStatusToHIPStatus(hcblasCreate(*handle)); + dummyGlobal = *handle; + return retVal; + +} + +inline static hipblasStatus_t hipblasDestroy(hipblasHandle_t& handle) { + return hipHCBLASStatusToHIPStatus(hcblasDestroy(handle)); +} + +//note: no handle +inline static hipblasStatus_t hipblasSetVector(int n, int elemSize, const void *x, int incx, void *y, int incy){ + return hipHCBLASStatusToHIPStatus(hcblasSetVector(dummyGlobal, n, elemSize, x, incx, y, incy)); //HGSOS no need for handle moving forward +} + +//note: no handle +inline static hipblasStatus_t hipblasGetVector(int n, int elemSize, const void *x, int incx, void *y, int incy){ + return hipHCBLASStatusToHIPStatus(hcblasGetVector(dummyGlobal, n, elemSize, x, incx, y, incy)); //HGSOS no need for handle +} + +//note: no handle +inline static hipblasStatus_t hipblasSetMatrix(int rows, int cols, int elemSize, const void *A, int lda, void *B, int ldb){ + return hipHCBLASStatusToHIPStatus(hcblasSetMatrix(dummyGlobal, rows, cols, elemSize, A, lda, B, ldb)); +} + +//note: no handle +inline static hipblasStatus_t hipblasGetMatrix(int rows, int cols, int elemSize, const void *A, int lda, void *B, int ldb){ + return hipHCBLASStatusToHIPStatus(hcblasGetMatrix(dummyGlobal, rows, cols, elemSize, A, lda, B, ldb)); +} + +inline static hipblasStatus_t hipblasSasum(hipblasHandle_t handle, int n, float *x, int incx, float *result){ + return hipHCBLASStatusToHIPStatus(hcblasSasum(handle, n, x, incx, result)); +} + +inline static hipblasStatus_t hipblasDasum(hipblasHandle_t handle, int n, double *x, int incx, double *result){ + return hipHCBLASStatusToHIPStatus(hcblasDasum(handle, n, x, incx, result)); +} + +inline static hipblasStatus_t hipblasSasumBatched(hipblasHandle_t handle, int n, float *x, int incx, float *result, int batchCount){ + return hipHCBLASStatusToHIPStatus(hcblasSasumBatched( handle, n, x, incx, result, batchCount)); +} + +inline static hipblasStatus_t hipblasDasumBatched(hipblasHandle_t handle, int n, double *x, int incx, double *result, int batchCount){ + return hipHCBLASStatusToHIPStatus(hcblasDasumBatched(handle, n, x, incx, result, batchCount)); +} + +inline static hipblasStatus_t hipblasSaxpy(hipblasHandle_t handle, int n, const float *alpha, const float *x, int incx, float *y, int incy) { + return hipHCBLASStatusToHIPStatus(hcblasSaxpy(handle, n, alpha, x, incx, y, incy)); +} + +inline static hipblasStatus_t hipblasSaxpyBatched(hipblasHandle_t handle, int n, const float *alpha, const float *x, int incx, float *y, int incy, int batchCount){ + return hipHCBLASStatusToHIPStatus(hcblasSaxpyBatched(handle, n, alpha, x, incx, y, incy, batchCount)); +} + +inline static hipblasStatus_t hipblasScopy(hipblasHandle_t handle, int n, const float *x, int incx, float *y, int incy){ + return hipHCBLASStatusToHIPStatus(hcblasScopy( handle, n, x, incx, y, incy)); +} + +inline static hipblasStatus_t hipblasDcopy(hipblasHandle_t handle, int n, const double *x, int incx, double *y, int incy){ + return hipHCBLASStatusToHIPStatus(hcblasDcopy( handle, n, x, incx, y, incy)); +} + +inline static hipblasStatus_t hipblasScopyBatched(hipblasHandle_t handle, int n, const float *x, int incx, float *y, int incy, int batchCount){ + return hipHCBLASStatusToHIPStatus(hcblasScopyBatched( handle, n, x, incx, y, incy, batchCount)); +} + +inline static hipblasStatus_t hipblasDcopyBatched(hipblasHandle_t handle, int n, const double *x, int incx, double *y, int incy, int batchCount){ + return hipHCBLASStatusToHIPStatus(hcblasDcopyBatched( handle, n, x, incx, y, incy, batchCount)); +} + +inline static hipblasStatus_t hipblasSdot (hipblasHandle_t handle, int n, const float *x, int incx, const float *y, int incy, float *result){ + return hipHCBLASStatusToHIPStatus(hcblasSdot(handle, n, x, incx, y, incy, result)); +} + +inline static hipblasStatus_t hipblasDdot (hipblasHandle_t handle, int n, const double *x, int incx, const double *y, int incy, double *result){ + return hipHCBLASStatusToHIPStatus(hcblasDdot(handle, n, x, incx, y, incy, result)); +} + +inline static hipblasStatus_t hipblasSdotBatched (hipblasHandle_t handle, int n, const float *x, int incx, const float *y, int incy, float *result, int batchCount){ + return hipHCBLASStatusToHIPStatus(hcblasSdotBatched(handle, n, x, incx, y, incy, result, batchCount)); +} + +inline static hipblasStatus_t hipblasDdotBatched (hipblasHandle_t handle, int n, const double *x, int incx, const double *y, int incy, double *result, int batchCount){ + return hipHCBLASStatusToHIPStatus(hcblasDdotBatched ( handle, n, x, incx, y, incy, result, batchCount)); +} + +inline static hipblasStatus_t hipblasSscal(hipblasHandle_t handle, int n, const float *alpha, float *x, int incx){ + return hipHCBLASStatusToHIPStatus(hcblasSscal(handle, n, alpha, x, incx)); +} + +inline static hipblasStatus_t hipblasDscal(hipblasHandle_t handle, int n, const double *alpha, double *x, int incx){ + return hipHCBLASStatusToHIPStatus(hcblasDscal(handle, n, alpha, x, incx)); +} + +inline static hipblasStatus_t hipblasSscalBatched(hipblasHandle_t handle, int n, const float *alpha, float *x, int incx, int batchCount){ + return hipHCBLASStatusToHIPStatus(hcblasSscalBatched(handle, n, alpha, x, incx, batchCount)); +} + +inline static hipblasStatus_t hipblasDscalBatched(hipblasHandle_t handle, int n, const double *alpha, double *x, int incx, int batchCount){ + return hipHCBLASStatusToHIPStatus(hcblasDscalBatched(handle, n, alpha, x, incx, batchCount)); +} + +inline static hipblasStatus_t hipblasSgemv(hipblasHandle_t handle, hipblasOperation_t trans, int m, int n, const float *alpha, float *A, int lda, + float *x, int incx, const float *beta, float *y, int incy){ + return hipHCBLASStatusToHIPStatus(hcblasSgemv(handle, hipOperationToHCCOperation(trans), m, n, alpha, A, lda, x, incx, beta, y, incy)); +} + +inline static hipblasStatus_t hipblasSgemvBatched(hipblasHandle_t handle, hipblasOperation_t trans, int m, int n, const float *alpha, float *A, int lda, + float *x, int incx, const float *beta, float *y, int incy, int batchCount){ + return hipHCBLASStatusToHIPStatus(hcblasSgemvBatched(handle, hipOperationToHCCOperation(trans), m, n, alpha, A, lda, x, incx, beta, y, incy, batchCount)); +} + +inline static hipblasStatus_t hipblasSger(hipblasHandle_t handle, int m, int n, const float *alpha, const float *x, int incx, const float *y, int incy, float *A, int lda){ + return hipHCBLASStatusToHIPStatus(hcblasSger(handle, m, n, alpha, x, incx, y, incy, A, lda)); +} + +inline static hipblasStatus_t hipblasSgerBatched(hipblasHandle_t handle, int m, int n, const float *alpha, const float *x, int incx, const float *y, int incy, float *A, int lda, int batchCount){ + return hipHCBLASStatusToHIPStatus(hcblasSgerBatched(handle, m, n, alpha, x, incx, y, incy, A, lda, batchCount)); +} + +inline static hipblasStatus_t hipblasSgemm(hipblasHandle_t handle, hipblasOperation_t transa, hipblasOperation_t transb, + int m, int n, int k, const float *alpha, float *A, int lda, float *B, int ldb, const float *beta, float *C, int ldc){ + return hipHCBLASStatusToHIPStatus(hcblasSgemm( handle, hipOperationToHCCOperation(transa), hipOperationToHCCOperation(transb), m, n, k, alpha, A, lda, B, ldb, beta, C, ldc)); +} + +inline static hipblasStatus_t hipblasCgemm(hipblasHandle_t handle, hipblasOperation_t transa, hipblasOperation_t transb, + int m, int n, int k, const hipComplex *alpha, hipComplex *A, int lda, hipComplex *B, int ldb, const hipComplex *beta, hipComplex *C, int ldc){ + return hipHCBLASStatusToHIPStatus(hcblasCgemm( handle, hipOperationToHCCOperation(transa), hipOperationToHCCOperation(transb), m, n, k, alpha, A, lda, B, ldb, beta, C, ldc)); +} + +inline static hipblasStatus_t hipblasSgemmBatched(hipblasHandle_t handle, hipblasOperation_t transa, hipblasOperation_t transb, + int m, int n, int k, const float *alpha, float *A, int lda, float *B, int ldb, const float *beta, float *C, int ldc, int batchCount){ + return hipHCBLASStatusToHIPStatus(hcblasSgemmBatched( handle, hipOperationToHCCOperation(transa), hipOperationToHCCOperation(transb), m, n, k, alpha, A, lda, B, ldb, beta, C, ldc, batchCount)); +} + +inline static hipblasStatus_t hipblasCgemmBatched(hipblasHandle_t handle, hipblasOperation_t transa, hipblasOperation_t transb, + int m, int n, int k, const hipComplex *alpha, hipComplex *A, int lda, hipComplex *B, int ldb, const hipComplex *beta, hipComplex *C, int ldc, int batchCount){ + return HIPBLAS_STATUS_NOT_SUPPORTED; + //return hipHCBLASStatusToHIPStatus(hcblasCgemmBatched( handle, transa, transb, m, n, k, alpha, A, lda, B, ldb, beta, C, ldc, batchCount)); +} + + + + +#ifdef __cplusplus +} +#endif + + diff --git a/projects/hip/include/hcc_detail/hip_hcc.h b/projects/hip/include/hcc_detail/hip_hcc.h index 405e2de67c..2f0c33a0d2 100644 --- a/projects/hip/include/hcc_detail/hip_hcc.h +++ b/projects/hip/include/hcc_detail/hip_hcc.h @@ -22,7 +22,7 @@ THE SOFTWARE. #include #include "hip/hcc_detail/hip_util.h" -#include "hip/hcc_detail/staging_buffer.h" +#include "hip/hcc_detail/unpinned_copy_engine.h" #if defined(__HCC__) && (__hcc_workweek__ < 16186) @@ -66,11 +66,16 @@ extern int HIP_VISIBLE_DEVICES; /* Contains a comma-separated sequence of GPU id extern int HIP_DISABLE_HW_KERNEL_DEP; extern int HIP_DISABLE_HW_COPY_DEP; -extern thread_local int tls_defaultDevice; +//--- +//Extern tls extern thread_local hipError_t tls_lastHipError; + + +//--- +//Forward defs: class ihipStream_t; class ihipDevice_t; - +class ihipCtx_t; // Color defs for debug messages: #define KNRM "\x1B[0m" @@ -90,13 +95,13 @@ class ihipDevice_t; #define STREAM_THREAD_SAFE 1 -#define DEVICE_THREAD_SAFE 1 +#define CTX_THREAD_SAFE 1 -// If FORCE_COPY_DEP=1 , HIP runtime will add +// If FORCE_COPY_DEP=1 , HIP runtime will add // synchronization for copy commands in the same stream, regardless of command type. // If FORCE_COPY_DEP=0 data copies of the same kind (H2H, H2D, D2H, D2D) are assumed to be implicitly ordered. -// ROCR runtime implementation currently provides this guarantee when using SDMA queues but not -// when using shader queues. +// ROCR runtime implementation currently provides this guarantee when using SDMA queues but not +// when using shader queues. // TODO - measure if this matters for performance, in particular for back-to-back small copies. // If not, we can simplify the copy dependency tracking by collapsing to a single Copy type, and always forcing dependencies for copy commands. #define FORCE_SAMEDIR_COPY_DEP 1 @@ -112,7 +117,7 @@ class ihipDevice_t; // 0x2 = prints a simple message with function name + return code when function exits. // 0x3 = print both. // Must be enabled at runtime with HIP_TRACE_API -#define COMPILE_HIP_TRACE_API 0x3 +#define COMPILE_HIP_TRACE_API 0x3 // Compile code that generates trace markers for CodeXL ATP at HIP function begin/end. @@ -132,13 +137,13 @@ class ihipDevice_t; #if COMPILE_HIP_ATP_MARKER #include "CXLActivityLogger.h" #define SCOPED_MARKER(markerName,group,userString) amdtScopedMarker(markerName, group, userString) -#else +#else // Swallow scoped markers: -#define SCOPED_MARKER(markerName,group,userString) +#define SCOPED_MARKER(markerName,group,userString) #endif -#if COMPILE_HIP_ATP_MARKER || (COMPILE_HIP_TRACE_API & 0x1) +#if COMPILE_HIP_ATP_MARKER || (COMPILE_HIP_TRACE_API & 0x1) #define API_TRACE(...)\ {\ if (HIP_ATP_MARKER || (COMPILE_HIP_DB && HIP_TRACE_API)) {\ @@ -163,15 +168,15 @@ class ihipDevice_t; std::call_once(hip_initialized, ihipInit);\ API_TRACE(__VA_ARGS__); -#define ihipLogStatus(_hip_status) \ +#define ihipLogStatus(hipStatus) \ ({\ - hipError_t _local_hip_status = _hip_status; /*local copy so _hip_status only evaluated once*/ \ - tls_lastHipError = _local_hip_status;\ + hipError_t localHipStatus = hipStatus; /*local copy so hipStatus only evaluated once*/ \ + tls_lastHipError = localHipStatus;\ \ if ((COMPILE_HIP_TRACE_API & 0x2) && HIP_TRACE_API) {\ - fprintf(stderr, " %ship-api: %-30s ret=%2d (%s)>>\n" KNRM, (_local_hip_status == 0) ? API_COLOR:KRED, __func__, _local_hip_status, ihipErrorString(_local_hip_status));\ + fprintf(stderr, " %ship-api: %-30s ret=%2d (%s)>>\n" KNRM, (localHipStatus == 0) ? API_COLOR:KRED, __func__, localHipStatus, ihipErrorString(localHipStatus));\ }\ - _local_hip_status;\ + localHipStatus;\ }) @@ -189,7 +194,7 @@ class ihipDevice_t; static const char *dbName [] = { - KNRM "hip-api", // not used, + KNRM "hip-api", // not used, KYEL "hip-sync", KCYN "hip-mem", KMAG "hip-copy1", @@ -205,17 +210,21 @@ static const char *dbName [] = fprintf (stderr, "%s", KNRM); \ }\ } -#else +#else /* Compile to empty code */ -#define tprintf(trace_level, ...) +#define tprintf(trace_level, ...) #endif + + + + class ihipException : public std::exception { public: ihipException(hipError_t e) : _code(e) {}; - hipError_t _code; + hipError_t _code; }; @@ -223,10 +232,6 @@ public: extern "C" { #endif -typedef class ihipStream_t* hipStream_t; -//typedef struct hipEvent_t { -// struct ihipEvent_t *_handle; -//} hipEvent_t; #ifdef __cplusplus } @@ -237,7 +242,7 @@ const hipStream_t hipStreamNull = 0x0; enum ihipCommand_t { ihipCommandCopyH2H, - ihipCommandCopyH2D, + ihipCommandCopyH2D, ihipCommandCopyD2H, ihipCommandCopyD2D, ihipCommandCopyP2P, @@ -258,9 +263,9 @@ typedef uint64_t SIGSEQNUM; // TODO-someday refactor this class so it can be stored in a vector<> // we already store the index here so we can use for garbage collection. struct ihipSignal_t { - hsa_signal_t _hsa_signal; // hsa signal handle + hsa_signal_t _hsaSignal; // hsa signal handle int _index; // Index in pool, used for garbage collection. - SIGSEQNUM _sig_id; // unique sequentially increasing ID. + SIGSEQNUM _sigId; // unique sequentially increasing ID. ihipSignal_t(); ~ihipSignal_t(); @@ -286,10 +291,11 @@ typedef std::mutex StreamMutex; typedef FakeMutex StreamMutex; #endif -#if DEVICE_THREAD_SAFE -typedef std::mutex DeviceMutex; +// Pair Device and Ctx together, these could also be toggled separately if desired. +#if CTX_THREAD_SAFE +typedef std::mutex CtxMutex; #else -typedef FakeMutex DeviceMutex; +typedef FakeMutex CtxMutex; #warning "Device thread-safe disabled" #endif @@ -301,7 +307,7 @@ template class LockedAccessor { public: - LockedAccessor(T &criticalData, bool autoUnlock=true) : + LockedAccessor(T &criticalData, bool autoUnlock=true) : _criticalData(&criticalData), _autoUnlock(autoUnlock) @@ -309,14 +315,14 @@ public: _criticalData->_mutex.lock(); }; - ~LockedAccessor() + ~LockedAccessor() { if (_autoUnlock) { _criticalData->_mutex.unlock(); } } - void unlock() + void unlock() { _criticalData->_mutex.unlock(); } @@ -333,7 +339,7 @@ private: template struct LockedBase { - // Experts-only interface for explicit locking. + // Experts-only interface for explicit locking. // Most uses should use the lock-accessor. void lock() { _mutex.lock(); } void unlock() { _mutex.unlock(); } @@ -342,8 +348,8 @@ struct LockedBase { }; -template -class ihipStreamCriticalBase_t : public LockedBase +template +class ihipStreamCriticalBase_t : public LockedBase { public: ihipStreamCriticalBase_t() : @@ -351,7 +357,7 @@ public: _last_copy_signal(NULL), _signalCursor(0), _oldest_live_sig_id(1), - _stream_sig_id(0), + _streamSigId(0), _kernelCnt(0), _signalCnt(0) { @@ -379,24 +385,22 @@ public: int _signalCursor; SIGSEQNUM _oldest_live_sig_id; // oldest live seq_id, anything < this can be allocated. std::deque _signalPool; // Pool of signals for use by this stream. - uint32_t _signalCnt; // Count of inflight commands using signals from the signal pool. - // Each copy may use 1-2 signals depending on command transitions: + uint32_t _signalCnt; // Count of inflight commands using signals from the signal pool. + // Each copy may use 1-2 signals depending on command transitions: // 2 are required if a barrier packet is inserted. uint32_t _kernelCnt; // Count of inflight kernels in this stream. Reset at ::wait(). - SIGSEQNUM _stream_sig_id; // Monotonically increasing unique signal id. + SIGSEQNUM _streamSigId; // Monotonically increasing unique signal id. }; -typedef ihipStreamCriticalBase_t ihipStreamCritical_t; +typedef ihipStreamCriticalBase_t ihipStreamCritical_t; typedef LockedAccessor LockedAccessor_StreamCrit_t; - - // Internal stream structure. class ihipStream_t { public: typedef uint64_t SeqNum_t ; - ihipStream_t(unsigned device_index, hc::accelerator_view av, unsigned int flags); + ihipStream_t(ihipCtx_t *ctx, hc::accelerator_view av, unsigned int flags); ~ihipStream_t(); // kind is hipMemcpyKind @@ -422,13 +426,14 @@ typedef uint64_t SeqNum_t ; // Non-threadsafe accessors - must be protected by high-level stream lock with accessor passed to function. - SIGSEQNUM lastCopySeqId (LockedAccessor_StreamCrit_t &crit) const { return crit->_last_copy_signal ? crit->_last_copy_signal->_sig_id : 0; }; + SIGSEQNUM lastCopySeqId (LockedAccessor_StreamCrit_t &crit) const { return crit->_last_copy_signal ? crit->_last_copy_signal->_sigId : 0; }; ihipSignal_t * allocSignal (LockedAccessor_StreamCrit_t &crit); //-- Non-racy accessors: // These functions access fields set at initialization time and are non-racy (so do not acquire mutex) - ihipDevice_t * getDevice() const; + const ihipDevice_t * getDevice() const; + ihipCtx_t * getCtx() const; public: @@ -439,45 +444,26 @@ public: unsigned _flags; private: - // Critical Data. THis MUST be accessed through LockedAccessor_StreamCrit_t - ihipStreamCritical_t _criticalData; - -private: - void enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal, ihipSignal_t *completionSignal); - void waitCopy(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *signal); - + void enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal, ihipSignal_t *completionSignal); + void waitCopy(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *signal); // The unsigned return is hipMemcpyKind unsigned resolveMemcpyDirection(bool srcTracked, bool dstTracked, bool srcInDeviceMem, bool dstInDeviceMem); - void setAsyncCopyAgents(unsigned kind, ihipCommand_t *commandType, hsa_agent_t *srcAgent, hsa_agent_t *dstAgent); + void setAsyncCopyAgents(unsigned kind, ihipCommand_t *commandType, hsa_agent_t *srcAgent, hsa_agent_t *dstAgent); - unsigned _device_index; // index into the g_device array +private: // Data + // Critical Data. THis MUST be accessed through LockedAccessor_StreamCrit_t + ihipStreamCritical_t _criticalData; + + ihipCtx_t *_ctx; // parent context that owns this stream. + + // Friends: friend std::ostream& operator<<(std::ostream& os, const ihipStream_t& s); }; -inline std::ostream& operator<<(std::ostream& os, const ihipStream_t& s) -{ - os << "stream#"; - os << s._device_index; - os << '.'; - os << s._id; - return os; -} - -inline std::ostream & operator<<(std::ostream& os, const dim3& s) -{ - os << '{'; - os << s.x; - os << ','; - os << s.y; - os << ','; - os << s.z; - os << '}'; - return os; -} //---- // Internal event structure: @@ -499,219 +485,199 @@ struct ihipEvent_t { hc::completion_future _marker; uint64_t _timestamp; // store timestamp, may be set on host or by marker. - SIGSEQNUM _copy_seq_id; + SIGSEQNUM _copySeqId; } ; -//--- -// Data that must be protected with thread-safe access -// All members are private - this class must be accessed through friend LockedAccessor which -// will lock the mutex on construction and unlock on destruction. -// -// MUTEX_TYPE is template argument so can easily convert to FakeMutex for performance or stress testing. -template -class ihipDeviceCriticalBase_t : LockedBase +//---- +// Properties of the HIP device. +// Multiple contexts can point to same device. +class ihipDevice_t { public: - ihipDeviceCriticalBase_t() : _stream_id(0), _peerAgents(nullptr) {}; - - void init(unsigned deviceCnt) { - assert(_peerAgents == nullptr); + ihipDevice_t(unsigned deviceId, unsigned deviceCnt, hc::accelerator &acc); + ~ihipDevice_t(); + + // Accessors: + ihipCtx_t *getPrimaryCtx() const { return _primaryCtx; }; + +public: + unsigned _deviceId; // device ID + + hc::accelerator _acc; + hsa_agent_t _hsaAgent; // hsa agent handle + + //! Number of compute units supported by the device: + unsigned _computeUnits; + hipDeviceProp_t _props; // saved device properties. + + UnpinnedCopyEngine *_stagingBuffer[2]; // one buffer for each direction. + int _isLargeBar; + + ihipCtx_t *_primaryCtx; + +private: + hipError_t initProperties(hipDeviceProp_t* prop); +}; +//============================================================================= + + + +//============================================================================= +//class ihipCtxCriticalBase_t +template +class ihipCtxCriticalBase_t : LockedBase +{ +public: + ihipCtxCriticalBase_t(unsigned deviceCnt) : + _peerCnt(0) + { _peerAgents = new hsa_agent_t[deviceCnt]; }; - ~ihipDeviceCriticalBase_t() { + ~ihipCtxCriticalBase_t() { if (_peerAgents != nullptr) { delete _peerAgents; _peerAgents = nullptr; } + _peerCnt = 0; } - friend class LockedAccessor; + // Streams: + void addStream(ihipStream_t *stream); std::list &streams() { return _streams; }; const std::list &const_streams() const { return _streams; }; - // "Allocate" a stream ID: - ihipStream_t::SeqNum_t incStreamId() { return _stream_id++; }; - bool isPeer(const ihipDevice_t *peer); // returns Trus if peer has access to memory physically located on this device. - bool addPeer(ihipDevice_t *peer); - bool removePeer(ihipDevice_t *peer); - void resetPeers(ihipDevice_t *thisDevice); - - - void addStream(ihipStream_t *stream); + // Peer Accessor classes: + bool isPeer(const ihipCtx_t *peer); // returns Trus if peer has access to memory physically located on this device. + bool addPeer(ihipCtx_t *peer); + bool removePeer(ihipCtx_t *peer); + void resetPeers(ihipCtx_t *thisDevice); uint32_t peerCnt() const { return _peerCnt; }; hsa_agent_t *peerAgents() const { return _peerAgents; }; -private: - //std::list< std::shared_ptr > _streams; // streams associated with this device. TODO - convert to shared_ptr. - std::list< ihipStream_t* > _streams; // streams associated with this device. - ihipStream_t::SeqNum_t _stream_id; + friend class LockedAccessor; +private: + //--- Stream Tracker: + std::list< ihipStream_t* > _streams; // streams associated with this device. + + + //--- Peer Tracker: // These reflect the currently Enabled set of peers for this GPU: // Enabled peers have permissions to access the memory physically allocated on this device. - std::list _peers; // list of enabled peer devices. + std::list _peers; // list of enabled peer devices. uint32_t _peerCnt; // number of enabled peers - hsa_agent_t *_peerAgents; // efficient packed array of enabled agents (to use for allocations.) + hsa_agent_t *_peerAgents; // efficient packed array of enabled agents (to use for allocations.) private: void recomputePeerAgents(); }; - -// Note Mutex selected based on DeviceMutex -typedef ihipDeviceCriticalBase_t ihipDeviceCritical_t; +// Note Mutex type Real/Fake selected based on CtxMutex +typedef ihipCtxCriticalBase_t ihipCtxCritical_t; // This type is used by functions that need access to the critical device structures. -typedef LockedAccessor LockedAccessor_DeviceCrit_t; +typedef LockedAccessor LockedAccessor_CtxCrit_t; +//============================================================================= - -//------------------------------------------------------------------------------------------------- -// Functions which read or write the critical data are named locked_. -// ihipDevice_t does not use recursive locks so the ihip implementation must avoid calling a locked_ function from within a locked_ function. -// External functions which call several locked_ functions will acquire and release the lock for each function. if this occurs in -// performance-sensitive code we may want to refactor by adding non-locked functions and creating a new locked_ member function to call them all. -class ihipDevice_t +//============================================================================= +//class ihipCtx_t: +// A HIP CTX (context) points at one of the existing devices and contains the streams, +// peer-to-peer mappings, creation flags. Multiple contexts can point to the same +// device. +// +class ihipCtx_t { public: // Functions: - ihipDevice_t() {}; // note: calls constructor for _criticalData - void init(unsigned device_index, unsigned deviceCnt, hc::accelerator &acc, unsigned flags); - ~ihipDevice_t(); + ihipCtx_t(ihipDevice_t *device, unsigned deviceCnt, unsigned flags); // note: calls constructor for _criticalData + ~ihipCtx_t(); + // Functions which read or write the critical data are named locked_. + // ihipCtx_t does not use recursive locks so the ihip implementation must avoid calling a locked_ function from within a locked_ function. + // External functions which call several locked_ functions will acquire and release the lock for each function. if this occurs in + // performance-sensitive code we may want to refactor by adding non-locked functions and creating a new locked_ member function to call them all. void locked_addStream(ihipStream_t *s); void locked_removeStream(ihipStream_t *s); void locked_reset(); void locked_waitAllStreams(); void locked_syncDefaultStream(bool waitOnSelf); - ihipDeviceCritical_t &criticalData() { return _criticalData; }; // TODO, move private. Fix P2P. + ihipCtxCritical_t &criticalData() { return _criticalData; }; // TODO, move private. Fix P2P. -public: // Data, set at initialization: - unsigned _device_index; // index into g_devices. + const ihipDevice_t *getDevice() const { return _device; }; - hipDeviceProp_t _props; // saved device properties. - hc::accelerator _acc; - hsa_agent_t _hsa_agent; // hsa agent handle + // TODO - review uses of getWriteableDevice(), can these be converted to getDevice() + ihipDevice_t *getWriteableDevice() const { return _device; }; +public: // Data // The NULL stream is used if no other stream is specified. - // NULL has special synchronization properties with other streams. - ihipStream_t *_default_stream; + // Default stream has special synchronization properties with other streams. + ihipStream_t *_defaultStream; - - unsigned _compute_units; - - StagingBuffer *_staging_buffer[2]; // one buffer for each direction. - int isLargeBar; - - unsigned _device_flags; + // Flags specified when the context is created: + unsigned _ctxFlags; private: - hipError_t getProperties(hipDeviceProp_t* prop); + ihipDevice_t *_device; + private: // Critical data, protected with locked access: // Members of _protected data MUST be accessed through the LockedAccessor. - // Search for LockedAccessor for examples; do not access _criticalData directly. - ihipDeviceCritical_t _criticalData; + // Search for LockedAccessor for examples; do not access _criticalData directly. + ihipCtxCritical_t _criticalData; }; +//================================================================================================= // Global variable definition: extern std::once_flag hip_initialized; -extern ihipDevice_t *g_devices; // Array of all non-emulated (ie GPU) accelerators in the system. -extern bool g_visible_device; // Set the flag when HIP_VISIBLE_DEVICES is set extern unsigned g_deviceCnt; -extern std::vector g_hip_visible_devices; /* vector of integers that contains the visible device IDs */ extern hsa_agent_t g_cpu_agent ; // the CPU agent. + //================================================================================================= -void ihipInit(); -const char *ihipErrorString(hipError_t); -ihipDevice_t *ihipGetTlsDefaultDevice(); -ihipDevice_t *ihipGetDevice(int); -void ihipSetTs(hipEvent_t e); +// Extern functions: +extern void ihipInit(); +extern const char *ihipErrorString(hipError_t); +extern ihipCtx_t *ihipGetTlsDefaultCtx(); +extern void ihipSetTlsDefaultCtx(ihipCtx_t *ctx); -template -hc::completion_future ihipMemcpyKernel(hipStream_t, T*, const T*, size_t); +extern ihipDevice_t *ihipGetDevice(int); +ihipCtx_t * ihipGetPrimaryCtx(unsigned deviceIndex); + +extern void ihipSetTs(hipEvent_t e); -template -hc::completion_future ihipMemsetKernel(hipStream_t, T*, T, size_t); hipStream_t ihipSyncAndResolveStream(hipStream_t); -template -hc::completion_future -ihipMemsetKernel(hipStream_t stream, T * ptr, T val, size_t sizeBytes) +// Stream printf functions: +inline std::ostream& operator<<(std::ostream& os, const ihipStream_t& s) { - int wg = std::min((unsigned)8, stream->getDevice()->_compute_units); - const int threads_per_wg = 256; - - int threads = wg * threads_per_wg; - if (threads > sizeBytes) { - threads = ((sizeBytes + threads_per_wg - 1) / threads_per_wg) * threads_per_wg; - } - - - hc::extent<1> ext(threads); - auto ext_tile = ext.tile(threads_per_wg); - - hc::completion_future cf = - hc::parallel_for_each( - stream->_av, - ext_tile, - [=] (hc::tiled_index<1> idx) - __attribute__((hc)) - { - int offset = amp_get_global_id(0); - // TODO-HCC - change to hc_get_local_size() - int stride = amp_get_local_size(0) * hc_get_num_groups(0) ; - - for (int i=offset; i_deviceId;; + os << '.'; + os << s._id; + return os; } -template -hc::completion_future -ihipMemcpyKernel(hipStream_t stream, T * c, const T * a, size_t sizeBytes) +inline std::ostream & operator<<(std::ostream& os, const dim3& s) { - int wg = std::min((unsigned)8, stream->getDevice()->_compute_units); - const int threads_per_wg = 256; - - int threads = wg * threads_per_wg; - if (threads > sizeBytes) { - threads = ((sizeBytes + threads_per_wg - 1) / threads_per_wg) * threads_per_wg; - } - - - hc::extent<1> ext(threads); - auto ext_tile = ext.tile(threads_per_wg); - - hc::completion_future cf = - hc::parallel_for_each( - stream->_av, - ext_tile, - [=] (hc::tiled_index<1> idx) - __attribute__((hc)) - { - int offset = amp_get_global_id(0); - // TODO-HCC - change to hc_get_local_size() - int stride = amp_get_local_size(0) * hc_get_num_groups(0) ; - - for (int i=offset; i // for getDeviceProp #include +enum { +HIP_SUCCESS = 0, +HIP_ERROR_INVALID_VALUE, +HIP_ERROR_NOT_INITIALIZED, +HIP_ERROR_LAUNCH_OUT_OF_RESOURCES +}; + typedef struct { // 32-bit Atomics unsigned hasGlobalInt32Atomics : 1; ///< 32-bit integer atomics for global memory. @@ -142,26 +149,67 @@ typedef struct hipPointerAttribute_t { // Also update the hipCUDAErrorTohipError function in NVCC path. typedef enum hipError_t { - hipSuccess = 0 ///< Successful completion. - ,hipErrorMemoryAllocation ///< Memory allocation error. - ,hipErrorLaunchOutOfResources ///< Out of resources error. - ,hipErrorInvalidValue ///< One or more of the parameters passed to the API call is NULL or not in an acceptable range. - ,hipErrorInvalidResourceHandle ///< Resource handle (hipEvent_t or hipStream_t) invalid. - ,hipErrorInvalidDevice ///< DeviceID must be in range 0...#compute-devices. - ,hipErrorInvalidMemcpyDirection ///< Invalid memory copy direction - ,hipErrorInvalidDevicePointer ///< Invalid Device Pointer - ,hipErrorInitializationError ///< TODO comment from hipErrorInitializationError + hipSuccess = 0, ///< Successful completion. + hipErrorOutOfMemory = 2, + hipErrorNotInitialized = 3, + hipErrorDeinitialized = 4, + hipErrorProfilerDisabled = 5, + hipErrorProfilerNotInitialized = 6, + hipErrorProfilerAlreadyStarted = 7, + hipErrorProfilerAlreadyStopped = 8, + hipErrorInvalidImage = 200, + hipErrorInvalidContext = 201, ///< Produced when input context is invalid. + hipErrorContextAlreadyCurrent = 202, + hipErrorMapFailed = 205, + hipErrorUnmapFailed = 206, + hipErrorArrayIsMapped = 207, + hipErrorAlreadyMapped = 208, + hipErrorNoBinaryForGpu = 209, + hipErrorAlreadyAcquired = 210, + hipErrorNotMapped = 211, + hipErrorNotMappedAsArray = 212, + hipErrorNotMappedAsPointer = 213, + hipErrorECCNotCorrectable = 214, + hipErrorUnsupportedLimit = 215, + hipErrorContextAlreadyInUse = 216, + hipErrorPeerAccessUnsupported = 217, + hipErrorInvalidKernelFile = 218, ///< In CUDA DRV, it is CUDA_ERROR_INVALID_PTX + hipErrorInvalidGraphicsContext = 219, + hipErrorInvalidSource = 300, + hipErrorFileNotFound = 301, + hipErrorSharedObjectSymbolNotFound = 302, + hipErrorSharedObjectInitFailed = 303, + hipErrorOperatingSystem = 304, + hipErrorInvalidHandle = 400, + hipErrorNotFound = 500, + hipErrorIllegalAddress = 700, - ,hipErrorNoDevice ///< Call to hipGetDeviceCount returned 0 devices - ,hipErrorNotReady ///< Indicates that asynchronous operations enqueued earlier are not ready. This is not actually an error, but is used to distinguish from hipSuccess (which indicates completion). APIs that return this error include hipEventQuery and hipStreamQuery. - ,hipErrorUnknown ///< Unknown error. - ,hipErrorPeerAccessNotEnabled ///< Peer access was never enabled from the current device. - ,hipErrorPeerAccessAlreadyEnabled ///< Peer access was already enabled from the current device. - ,hipErrorRuntimeMemory ///< HSA runtime memory call returned error. Typically not seen in production systems. - ,hipErrorRuntimeOther ///< HSA runtime call other than memory returned error. Typically not seen in production systems. - ,hipErrorHostMemoryAlreadyRegistered ///< Produced when trying to lock a page-locked memory. - ,hipErrorHostMemoryNotRegistered ///< Produced when trying to unlock a non-page-locked memory. - ,hipErrorTbd ///< Marker that more error codes are needed. +// Runtime Error Codes start here. + hipErrorMissingConfiguration = 1, + hipErrorMemoryAllocation = 2, ///< Memory allocation error. + hipErrorInitializationError = 3, ///< TODO comment from hipErrorInitializationError + hipErrorLaunchFailure = 4, + hipErrorPriorLaunchFailure = 5, + hipErrorLaunchTimeOut = 6, + hipErrorLaunchOutOfResources = 7, ///< Out of resources error. + hipErrorInvalidDeviceFunction = 8, + hipErrorInvalidConfiguration = 9, + hipErrorInvalidDevice = 10, ///< DeviceID must be in range 0...#compute-devices. + hipErrorInvalidValue = 11, ///< One or more of the parameters passed to the API call is NULL or not in an acceptable range. + hipErrorInvalidDevicePointer = 17, ///< Invalid Device Pointer + hipErrorInvalidMemcpyDirection = 21, ///< Invalid memory copy direction + hipErrorUnknown = 30, ///< Unknown error. + hipErrorInvalidResourceHandle = 33, ///< Resource handle (hipEvent_t or hipStream_t) invalid. + hipErrorNotReady = 34, ///< Indicates that asynchronous operations enqueued earlier are not ready. This is not actually an error, but is used to distinguish from hipSuccess (which indicates completion). APIs that return this error include hipEventQuery and hipStreamQuery. + hipErrorNoDevice = 38, ///< Call to hipGetDeviceCount returned 0 devices + hipErrorPeerAccessAlreadyEnabled = 50, ///< Peer access was already enabled from the current device. + + hipErrorPeerAccessNotEnabled = 51, ///< Peer access was never enabled from the current device. + hipErrorRuntimeMemory, ///< HSA runtime memory call returned error. Typically not seen in production systems. + hipErrorRuntimeOther, ///< HSA runtime call other than memory returned error. Typically not seen in production systems. + hipErrorHostMemoryAlreadyRegistered = 61, ///< Produced when trying to lock a page-locked memory. + hipErrorHostMemoryNotRegistered = 62, ///< Produced when trying to unlock a non-page-locked memory. + hipErrorTbd ///< Marker that more error codes are needed. } hipError_t; /* diff --git a/projects/hip/include/hipblas.h b/projects/hip/include/hipblas.h new file mode 100644 index 0000000000..0e9b493a41 --- /dev/null +++ b/projects/hip/include/hipblas.h @@ -0,0 +1,66 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +//! HIP = Heterogeneous-compute Interface for Portability +//! +//! Define a extremely thin runtime layer that allows source code to be compiled unmodified +//! through either AMD HCC or NVCC. Key features tend to be in the spirit +//! and terminology of CUDA, but with a portable path to other accelerators as well. +//! +//! This is the master include file for hipblas, wrapping around hcblas and cublas "version 1" +// + +#pragma once + +enum hipblasStatus_t { + HIPBLAS_STATUS_SUCCESS, // Function succeeds + HIPBLAS_STATUS_NOT_INITIALIZED, // HIPBLAS library not initialized + HIPBLAS_STATUS_ALLOC_FAILED, // resource allocation failed + HIPBLAS_STATUS_INVALID_VALUE, // unsupported numerical value was passed to function + HIPBLAS_STATUS_MAPPING_ERROR, // access to GPU memory space failed + HIPBLAS_STATUS_EXECUTION_FAILED, // GPU program failed to execute + HIPBLAS_STATUS_INTERNAL_ERROR, // an internal HIPBLAS operation failed + HIPBLAS_STATUS_NOT_SUPPORTED // cublas supports this, but not hcblas +}; + +enum hipblasOperation_t { + HIPBLAS_OP_N, + HIPBLAS_OP_T, + HIPBLAS_OP_C +}; + +// Some standard header files, these are included by hc.hpp and so want to make them avail on both +// paths to provide a consistent include env and avoid "missing symbol" errors that only appears +// on NVCC path: + +#if defined(__HIP_PLATFORM_HCC__) and not defined (__HIP_PLATFORM_NVCC__) +#include +#elif defined(__HIP_PLATFORM_NVCC__) and not defined (__HIP_PLATFORM_HCC__) +#include +#else +#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__"); +#endif + + + + + diff --git a/projects/hip/include/nvcc_detail/hip_blas.h b/projects/hip/include/nvcc_detail/hip_blas.h new file mode 100644 index 0000000000..f01fb171c7 --- /dev/null +++ b/projects/hip/include/nvcc_detail/hip_blas.h @@ -0,0 +1,268 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#pragma once + +#include +#include +#include + +//HGSOS for Kalmar leave it as C++, only cublas needs C linkage. + +#ifdef __cplusplus +extern "C" { +#endif + + +typedef cublasHandle_t hipblasHandle_t ; +typedef cuComplex hipComplex; + +/* Unsupported types + "cublasFillMode_t", + "cublasDiagType_t", + "cublasSideMode_t", + "cublasPointerMode_t", + "cublasAtomicsMode_t", + "cublasDataType_t" +*/ + + +inline static cublasOperation_t hipOperationToCudaOperation( hipblasOperation_t op) +{ + switch (op) + { + case HIPBLAS_OP_N: + return CUBLAS_OP_N; + + case HIPBLAS_OP_T: + return CUBLAS_OP_T; + + case HIPBLAS_OP_C: + return CUBLAS_OP_C; + + default: + throw "Non existent OP"; + } +} + +inline static hipblasOperation_t CudaOperationToHIPOperation( cublasOperation_t op) +{ + switch (op) + { + case CUBLAS_OP_N : + return HIPBLAS_OP_N; + + case CUBLAS_OP_T : + return HIPBLAS_OP_T; + + case CUBLAS_OP_C : + return HIPBLAS_OP_C; + + default: + throw "Non existent OP"; + } +} + + +inline static hipblasStatus_t hipCUBLASStatusToHIPStatus(cublasStatus_t cuStatus) +{ + switch(cuStatus) + { + case CUBLAS_STATUS_SUCCESS: + return HIPBLAS_STATUS_SUCCESS; + case CUBLAS_STATUS_NOT_INITIALIZED: + return HIPBLAS_STATUS_NOT_INITIALIZED; + case CUBLAS_STATUS_ALLOC_FAILED: + return HIPBLAS_STATUS_ALLOC_FAILED; + case CUBLAS_STATUS_INVALID_VALUE: + return HIPBLAS_STATUS_INVALID_VALUE; + case CUBLAS_STATUS_MAPPING_ERROR: + return HIPBLAS_STATUS_MAPPING_ERROR; + case CUBLAS_STATUS_EXECUTION_FAILED: + return HIPBLAS_STATUS_EXECUTION_FAILED; + case CUBLAS_STATUS_INTERNAL_ERROR: + return HIPBLAS_STATUS_INTERNAL_ERROR; + case CUBLAS_STATUS_NOT_SUPPORTED: + return HIPBLAS_STATUS_NOT_SUPPORTED; + default: + throw "Unimplemented status"; + } +} + + +inline static hipblasStatus_t hipblasCreate(hipblasHandle_t* handle) { + return hipCUBLASStatusToHIPStatus(cublasCreate(&*handle)); +} + +//TODO broke common API semantics, think about this again. +inline static hipblasStatus_t hipblasDestroy(hipblasHandle_t handle) { + return hipCUBLASStatusToHIPStatus(cublasDestroy(handle)); +} + +//note: no handle +inline static hipblasStatus_t hipblasSetVector(int n, int elemSize, const void *x, int incx, void *y, int incy){ + return hipCUBLASStatusToHIPStatus(cublasSetVector(n, elemSize, x, incx, y, incy)); //HGSOS no need for handle +} + +//note: no handle +inline static hipblasStatus_t hipblasGetVector(int n, int elemSize, const void *x, int incx, void *y, int incy){ + return hipCUBLASStatusToHIPStatus(cublasGetVector(n, elemSize, x, incx, y, incy)); //HGSOS no need for handle +} + +//note: no handle +inline static hipblasStatus_t hipblasSetMatrix(int rows, int cols, int elemSize, const void *A, int lda, void *B, int ldb){ + return hipCUBLASStatusToHIPStatus(cublasSetMatrix(rows, cols, elemSize, A, lda, B, ldb)); +} + +//note: no handle +inline static hipblasStatus_t hipblasGetMatrix(int rows, int cols, int elemSize, const void *A, int lda, void *B, int ldb){ + return hipCUBLASStatusToHIPStatus(cublasGetMatrix(rows, cols, elemSize, A, lda, B, ldb)); +} + +inline static hipblasStatus_t hipblasSasum(hipblasHandle_t handle, int n, float *x, int incx, float *result){ + return hipCUBLASStatusToHIPStatus(cublasSasum(handle, n, x, incx, result)); +} + +inline static hipblasStatus_t hipblasDasum(hipblasHandle_t handle, int n, double *x, int incx, double *result){ + return hipCUBLASStatusToHIPStatus(cublasDasum( handle, n, x, incx, result)); +} + +inline static hipblasStatus_t hipblasSasumBatched(hipblasHandle_t handle, int n, float *x, int incx, float *result, int batchCount){ + //TODO warn user that function was demoted to ignore batch + return hipCUBLASStatusToHIPStatus(cublasSasum( handle, n, x, incx, result)); +} + +inline static hipblasStatus_t hipblasDasumBatched(hipblasHandle_t handle, int n, double *x, int incx, double *result, int batchCount){ + //TODO warn user that function was demoted to ignore batch + return hipCUBLASStatusToHIPStatus(cublasDasum(handle, n, x, incx, result)); +} + +inline static hipblasStatus_t hipblasSaxpy(hipblasHandle_t handle, int n, const float *alpha, const float *x, int incx, float *y, int incy) { + return hipCUBLASStatusToHIPStatus(cublasSaxpy(handle, n, alpha, x, incx, y, incy)); +} + +inline static hipblasStatus_t hipblasSaxpyBatched(hipblasHandle_t handle, int n, const float *alpha, const float *x, int incx, float *y, int incy, int batchCount){ + //TODO warn user that function was demoted to ignore batch + return hipCUBLASStatusToHIPStatus(cublasSaxpy(handle, n, alpha, x, incx, y, incy)); +} + +inline static hipblasStatus_t hipblasScopy(hipblasHandle_t handle, int n, const float *x, int incx, float *y, int incy){ + return hipCUBLASStatusToHIPStatus(cublasScopy( handle, n, x, incx, y, incy)); +} + +inline static hipblasStatus_t hipblasDcopy(hipblasHandle_t handle, int n, const double *x, int incx, double *y, int incy){ + return hipCUBLASStatusToHIPStatus(cublasDcopy( handle, n, x, incx, y, incy)); +} + +inline static hipblasStatus_t hipblasScopyBatched(hipblasHandle_t handle, int n, const float *x, int incx, float *y, int incy, int batchCount){ + //TODO warn user that function was demoted to ignore batch + return hipCUBLASStatusToHIPStatus(cublasScopy( handle, n, x, incx, y, incy)); +} + +inline static hipblasStatus_t hipblasDcopyBatched(hipblasHandle_t handle, int n, const double *x, int incx, double *y, int incy, int batchCount){ + //TODO warn user that function was demoted to ignore batch + return hipCUBLASStatusToHIPStatus(cublasDcopy( handle, n, x, incx, y, incy)); +} + +inline static hipblasStatus_t hipblasSdot (hipblasHandle_t handle, int n, const float *x, int incx, const float *y, int incy, float *result){ + return hipCUBLASStatusToHIPStatus(cublasSdot ( handle, n, x, incx, y, incy, result)); +} + +inline static hipblasStatus_t hipblasDdot (hipblasHandle_t handle, int n, const double *x, int incx, const double *y, int incy, double *result){ + return hipCUBLASStatusToHIPStatus(cublasDdot ( handle, n, x, incx, y, incy, result)); +} + +inline static hipblasStatus_t hipblasSdotBatched (hipblasHandle_t handle, int n, const float *x, int incx, const float *y, int incy, float *result, int batchCount){ + //TODO warn user that function was demoted to ignore batch + return hipCUBLASStatusToHIPStatus(cublasSdot ( handle, n, x, incx, y, incy, result)); +} + +inline static hipblasStatus_t hipblasDdotBatched (hipblasHandle_t handle, int n, const double *x, int incx, const double *y, int incy, double *result, int batchCount){ + //TODO warn user that function was demoted to ignore batch + return hipCUBLASStatusToHIPStatus(cublasDdot ( handle, n, x, incx, y, incy, result)); +} + +inline static hipblasStatus_t hipblasSscal(hipblasHandle_t handle, int n, const float *alpha, float *x, int incx){ + return hipCUBLASStatusToHIPStatus(cublasSscal(handle, n, alpha, x, incx)); +} +inline static hipblasStatus_t hipblasDscal(hipblasHandle_t handle, int n, const double *alpha, double *x, int incx){ + return hipCUBLASStatusToHIPStatus(cublasDscal(handle, n, alpha, x, incx)); +} +inline static hipblasStatus_t hipblasSscalBatched(hipblasHandle_t handle, int n, const float *alpha, float *x, int incx, int batchCount){ + //TODO warn user that function was demoted to ignore batch + return hipCUBLASStatusToHIPStatus(cublasSscal(handle, n, alpha, x, incx)); +} +inline static hipblasStatus_t hipblasDscalBatched(hipblasHandle_t handle, int n, const double *alpha, double *x, int incx, int batchCount){ + //TODO warn user that function was demoted to ignore batch + return hipCUBLASStatusToHIPStatus(cublasDscal(handle, n, alpha, x, incx)); +} + +inline static hipblasStatus_t hipblasSgemv(hipblasHandle_t handle, hipblasOperation_t trans, int m, int n, const float *alpha, float *A, int lda, + float *x, int incx, const float *beta, float *y, int incy){ + return hipCUBLASStatusToHIPStatus(cublasSgemv(handle, hipOperationToCudaOperation(trans), m, n, alpha, A, lda, x, incx, beta, y, incy)); +} + +inline static hipblasStatus_t hipblasSgemvBatched(hipblasHandle_t handle, hipblasOperation_t trans, int m, int n, const float *alpha, float *A, int lda, + float *x, int incx, const float *beta, float *y, int incy, int batchCount){ + //TODO warn user that function was demoted to ignore batch + return hipCUBLASStatusToHIPStatus(cublasSgemv(handle, hipOperationToCudaOperation(trans), m, n, alpha, A, lda, x, incx, beta, y, incy)); +} + +inline static hipblasStatus_t hipblasSger(hipblasHandle_t handle, int m, int n, const float *alpha, const float *x, int incx, const float *y, int incy, float *A, int lda){ + return hipCUBLASStatusToHIPStatus(cublasSger(handle, m, n, alpha, x, incx, y, incy, A, lda)); +} + +inline static hipblasStatus_t hipblasSgerBatched(hipblasHandle_t handle, int m, int n, const float *alpha, const float *x, int incx, const float *y, int incy, float *A, int lda, int batchCount){ + //TODO warn user that function was demoted to ignore batch + return hipCUBLASStatusToHIPStatus(cublasSger(handle, m, n, alpha, x, incx, y, incy, A, lda)); +} + +inline static hipblasStatus_t hipblasSgemm(hipblasHandle_t handle, hipblasOperation_t transa, hipblasOperation_t transb, + int m, int n, int k, const float *alpha, float *A, int lda, float *B, int ldb, const float *beta, float *C, int ldc){ + return hipCUBLASStatusToHIPStatus(cublasSgemm( handle, hipOperationToCudaOperation(transa), hipOperationToCudaOperation(transb), m, n, k, alpha, A, lda, B, ldb, beta, C, ldc)); +} + +inline static hipblasStatus_t hipblasCgemm(hipblasHandle_t handle, hipblasOperation_t transa, hipblasOperation_t transb, + int m, int n, int k, const hipComplex *alpha, hipComplex *A, int lda, hipComplex *B, int ldb, const hipComplex *beta, hipComplex *C, int ldc){ + return hipCUBLASStatusToHIPStatus(cublasCgemm( handle, hipOperationToCudaOperation(transa), hipOperationToCudaOperation(transb), m, n, k, alpha, A, lda, B, ldb, beta, C, ldc)); +} + +inline static hipblasStatus_t hipblasSgemmBatched(hipblasHandle_t handle, hipblasOperation_t transa, hipblasOperation_t transb, + int m, int n, int k, const float *alpha, float *A, int lda, float *B, int ldb, const float *beta, float *C, int ldc, int batchCount){ + //TODO incompatible API + return HIPBLAS_STATUS_NOT_SUPPORTED; + //return hipCUBLASStatusToHIPStatus(cublasSgemmBatched( handle, hipOperationToCudaOperation(transa), hipOperationToCudaOperation(transb), m, n, k, alpha, A, lda, B, ldb, beta, C, ldc, batchCount)); +} + + +inline static hipblasStatus_t hipblasCgemmBatched(hipblasHandle_t handle, hipblasOperation_t transa, hipblasOperation_t transb, + int m, int n, int k, const hipComplex *alpha, hipComplex *A, int lda, hipComplex *B, int ldb, const hipComplex *beta, hipComplex *C, int ldc, int batchCount){ + + //TODO incompatible API + return HIPBLAS_STATUS_NOT_SUPPORTED; + //return hipCUBLASStatusToHIPStatus(cublasCgemmBatched( handle, hipOperationToCudaOperation(transa), hipOperationToCudaOperation(transb), m, n, k, alpha, A, lda, B, ldb, beta, C, ldc, batchCount)); +} + +#ifdef __cplusplus +} +#endif + + diff --git a/projects/hip/include/nvcc_detail/hip_runtime.h b/projects/hip/include/nvcc_detail/hip_runtime.h index 9e05974b92..569d6297bf 100644 --- a/projects/hip/include/nvcc_detail/hip_runtime.h +++ b/projects/hip/include/nvcc_detail/hip_runtime.h @@ -99,6 +99,8 @@ kernelName<<>>(0, ##__VA_ARGS__);\ #define HIP_DYNAMIC_SHARED(type, var) \ extern __shared__ type var[]; \ +#define HIP_DYNAMIC_SHARED_ATTRIBUTE + #endif diff --git a/projects/hip/include/nvcc_detail/hip_runtime_api.h b/projects/hip/include/nvcc_detail/hip_runtime_api.h index 6ea49fb3a1..a51bca2d02 100644 --- a/projects/hip/include/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/nvcc_detail/hip_runtime_api.h @@ -82,7 +82,7 @@ switch(cuError) { case cudaErrorHostMemoryAlreadyRegistered : return hipErrorHostMemoryAlreadyRegistered ; case cudaErrorHostMemoryNotRegistered : return hipErrorHostMemoryNotRegistered ; default : return hipErrorUnknown; // Note - translated error. -}; +}; } // TODO match the error enum names of hip and cuda @@ -108,7 +108,7 @@ switch(hError) { case hipErrorHostMemoryNotRegistered : return cudaErrorHostMemoryNotRegistered ; case hipErrorTbd : return cudaErrorUnknown; // Note - translated error. default : return cudaErrorUnknown; // Note - translated error. -} +} } inline static cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind kind) { @@ -347,6 +347,31 @@ inline static hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t att return hipCUDAErrorTohipError(cerror); } +template +inline static hipError_t hipOccupancyMaxPotentialBlockSize( + int *minGridSize, + int *blockSize, + T func, + size_t dynamicSMemSize = 0, + int blockSizeLimit = 0, + unsigned int flags = 0 + ){ + cudaError_t cerror; + cerror = cudaOccupancyMaxPotentialBlockSize(minGridSize, blockSize, func, dynamicSMemSize, blockSizeLimit, flags); + return hipCUDAErrorTohipError(cerror); +} + +inline static hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor( + int *numBlocks, + const void* func, + int blockSize, + size_t dynamicSMemSize + ) +{ + cudaError_t cerror; + cerror = cudaOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, func, blockSize, dynamicSMemSize); + return hipCUDAErrorTohipError(cerror); +} inline static hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr){ cudaPointerAttributes cPA; diff --git a/projects/hip/packaging/hip_base.txt b/projects/hip/packaging/hip_base.txt index 1edba53901..67df3c10f6 100644 --- a/projects/hip/packaging/hip_base.txt +++ b/projects/hip/packaging/hip_base.txt @@ -3,6 +3,7 @@ project(hip_base) install(DIRECTORY @hip_SOURCE_DIR@/bin DESTINATION . USE_SOURCE_PERMISSIONS) install(DIRECTORY @hip_SOURCE_DIR@/include DESTINATION . PATTERN "hip" EXCLUDE) +install(FILES @PROJECT_BINARY_DIR@/.version DESTINATION bin) ############################# # Packaging steps diff --git a/projects/hip/samples/7_Advanced/hipblas_saxpy/Makefile b/projects/hip/samples/7_Advanced/hipblas_saxpy/Makefile new file mode 100644 index 0000000000..ed88be2dd0 --- /dev/null +++ b/projects/hip/samples/7_Advanced/hipblas_saxpy/Makefile @@ -0,0 +1,34 @@ +HIP_PATH?= $(wildcard /opt/rocm/hip) +ifeq (,$(HIP_PATH)) + HIP_PATH=../../.. +endif +HIPCC=$(HIP_PATH)/bin/hipcc + +HIPCC_FLAGS += -std=c++11 +HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform) +ifeq (${HIP_PLATFORM}, nvcc) + LIBS = -lcublas +endif +ifeq (${HIP_PLATFORM}, hcc) + HCBLAS_ROOT?= $(wildcard /opt/rocm/hcblas) + HIPCC_FLAGS += -stdlib=libc++ -I$(HCBLAS_ROOT)/include + LIBS = -L$(HCBLAS_ROOT)/lib -lhcblas +endif + + +all: saxpy.hipblas.out + +saxpy.cublas.out : saxpy.cublas.cpp + nvcc -std=c++11 -I$(CUDA_HOME)/include saxpy.cublas.cpp -o $@ -L$(CUDA_HOME)/lib64 -lcublas + +# $HIPBLAS_ROOT/bin/hipifyblas ./saxpy.cublas.cpp > ./saxpy.hipblas.cpp +# Then review & finish port in saxpy.hipblas.cpp + +saxpy.hipblasref.o: saxpy.hipblasref.cpp + $(HIPCC) $(HIPCC_FLAGS) -c $< -o $@ + +saxpy.hipblas.out: saxpy.hipblasref.o + $(HIPCC) $< -o $@ $(LIBS) + +clean: + rm -f *.o *.out diff --git a/projects/hip/samples/7_Advanced/hipblas_saxpy/saxpy.cublas.cpp b/projects/hip/samples/7_Advanced/hipblas_saxpy/saxpy.cublas.cpp new file mode 100644 index 0000000000..03a38f3fb1 --- /dev/null +++ b/projects/hip/samples/7_Advanced/hipblas_saxpy/saxpy.cublas.cpp @@ -0,0 +1,94 @@ + +#include +#include +#include +#include + +// header file for the GPU API +#include +#include + +#define N (1024 * 500) + +#define CHECK(cmd) \ +{\ + cudaError_t error = cmd; \ + if (error != cudaSuccess) { \ + fprintf(stderr, "error: '%s'(%d) at %s:%d\n", cudaGetErrorString(error), error,__FILE__, __LINE__); \ + exit(EXIT_FAILURE);\ + }\ +} + +#define CHECK_BLAS(cmd) \ +{\ + cublasStatus_t error = cmd;\ + if (error != CUBLAS_STATUS_SUCCESS) { \ + fprintf(stderr, "error: (%d) at %s:%d\n", error,__FILE__, __LINE__); \ + exit(EXIT_FAILURE);\ + }\ +} + +int main() { + + const float a = 100.0f; + float x[N]; + float y[N], y_cpu_res[N], y_gpu_res[N]; + + // initialize the input data + std::default_random_engine random_gen; + std::uniform_real_distribution distribution(-N, N); + std::generate_n(x, N, [&]() { return distribution(random_gen); }); + std::generate_n(y, N, [&]() { return distribution(random_gen); }); + std::copy_n(y, N, y_cpu_res); + + // Explicit GPU code: + + size_t Nbytes = N*sizeof(float); + float *x_gpu, *y_gpu; + + cublasHandle_t handle; + + cudaDeviceProp props; + CHECK(cudaGetDeviceProperties(&props, 0/*deviceID*/)); + printf ("info: running on device %s\n", props.name); + + printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0); + printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0); + CHECK(cudaMalloc(&x_gpu, Nbytes)); + CHECK(cudaMalloc(&y_gpu, Nbytes)); + + // Initialize the blas library + CHECK_BLAS ( cublasCreate(&handle)); + + // copy n elements from a vector in host memory space to a vector in GPU memory space + printf ("info: copy Host2Device\n"); + CHECK_BLAS ( cublasSetVector(N, sizeof(*x), x, 1, x_gpu, 1)); + CHECK_BLAS ( cublasSetVector(N, sizeof(*y), y, 1, y_gpu, 1)); + + printf ("info: launch 'saxpy' kernel\n"); + CHECK_BLAS ( cublasSaxpy(handle, N, &a, x_gpu, 1, y_gpu, 1)); + + cudaDeviceSynchronize(); + + printf ("info: copy Device2Host\n"); + CHECK_BLAS ( cublasGetVector(N, sizeof(*y_gpu_res), y_gpu, 1, y_gpu_res, 1)); + + // CPU implementation of saxpy + for (int i = 0; i < N; i++) { + y_cpu_res[i] = a * x[i] + y[i]; + } + + // verify the results + int errors = 0; + for (int i = 0; i < N; i++) { + if (fabs(y_cpu_res[i] - y_gpu_res[i]) > fabs(y_cpu_res[i] * 0.0001f)) + errors++; + } + std::cout << errors << " errors" << std::endl; + + CHECK( cudaFree(x_gpu)); + CHECK( cudaFree(y_gpu)); + CHECK_BLAS( cublasDestroy(handle)); + + return errors; +} diff --git a/projects/hip/samples/7_Advanced/hipblas_saxpy/saxpy.hipblasref.cpp b/projects/hip/samples/7_Advanced/hipblas_saxpy/saxpy.hipblasref.cpp new file mode 100644 index 0000000000..3f20c8a7cc --- /dev/null +++ b/projects/hip/samples/7_Advanced/hipblas_saxpy/saxpy.hipblasref.cpp @@ -0,0 +1,94 @@ + +#include +#include +#include +#include + +// header file for the GPU API +#include +#include + +#define N (1024 * 500) + +#define CHECK(cmd) \ +{\ + hipError_t error = cmd; \ + if (error != hipSuccess) { \ + fprintf(stderr, "error: '%s'(%d) at %s:%d\n", hipGetErrorString(error), error,__FILE__, __LINE__); \ + exit(EXIT_FAILURE);\ + }\ +} + +#define CHECK_BLAS(cmd) \ +{\ + hipblasStatus_t error = cmd;\ + if (error != HIPBLAS_STATUS_SUCCESS) { \ + fprintf(stderr, "error: (%d) at %s:%d\n", error,__FILE__, __LINE__); \ + exit(EXIT_FAILURE);\ + }\ +} + +int main() { + + const float a = 100.0f; + float x[N]; + float y[N], y_cpu_res[N], y_gpu_res[N]; + + // initialize the input data + std::default_random_engine random_gen; + std::uniform_real_distribution distribution(-N, N); + std::generate_n(x, N, [&]() { return distribution(random_gen); }); + std::generate_n(y, N, [&]() { return distribution(random_gen); }); + std::copy_n(y, N, y_cpu_res); + + // Explicit GPU code: + + size_t Nbytes = N*sizeof(float); + float *x_gpu, *y_gpu; + + hipblasHandle_t handle; + + hipDeviceProp_t props; + CHECK(hipGetDeviceProperties(&props, 0/*deviceID*/)); + printf ("info: running on device %s\n", props.name); + + printf ("info: allocate host mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0); + printf ("info: allocate device mem (%6.2f MB)\n", 2*Nbytes/1024.0/1024.0); + CHECK(hipMalloc(&x_gpu, Nbytes)); + CHECK(hipMalloc(&y_gpu, Nbytes)); + + // Initialize the blas library + CHECK_BLAS ( hipblasCreate(&handle)); + + // copy n elements from a vector in host memory space to a vector in GPU memory space + printf ("info: copy Host2Device\n"); + CHECK_BLAS ( hipblasSetVector(N, sizeof(*x), x, 1, x_gpu, 1)); + CHECK_BLAS ( hipblasSetVector(N, sizeof(*y), y, 1, y_gpu, 1)); + + printf ("info: launch 'saxpy' kernel\n"); + CHECK_BLAS ( hipblasSaxpy(handle, N, &a, x_gpu, 1, y_gpu, 1)); + + hipDeviceSynchronize(); + + printf ("info: copy Device2Host\n"); + CHECK_BLAS ( hipblasGetVector(N, sizeof(*y_gpu_res), y_gpu, 1, y_gpu_res, 1)); + + // CPU implementation of saxpy + for (int i = 0; i < N; i++) { + y_cpu_res[i] = a * x[i] + y[i]; + } + + // verify the results + int errors = 0; + for (int i = 0; i < N; i++) { + if (fabs(y_cpu_res[i] - y_gpu_res[i]) > fabs(y_cpu_res[i] * 0.0001f)) + errors++; + } + std::cout << errors << " errors" << std::endl; + + CHECK( hipFree(x_gpu)); + CHECK( hipFree(y_gpu)); + CHECK_BLAS( hipblasDestroy(handle)); + + return errors; +} diff --git a/projects/hip/src/device_util.cpp b/projects/hip/src/device_util.cpp index ca9d858981..ce829b73d3 100644 --- a/projects/hip/src/device_util.cpp +++ b/projects/hip/src/device_util.cpp @@ -1389,7 +1389,7 @@ __device__ double norm4d(double a, double b, double c, double d) double y = c*c + d*d; return hc::precise_math::sqrt(x+y); } -__device__ double normcdf(float y) +__device__ double normcdf(double y) { return ((hc::precise_math::erf(y)/HIP_SQRT_2) + 1)/2; } diff --git a/projects/hip/src/hip_context.cpp b/projects/hip/src/hip_context.cpp new file mode 100644 index 0000000000..ee9e37a1a1 --- /dev/null +++ b/projects/hip/src/hip_context.cpp @@ -0,0 +1,219 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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. +*/ + +//--- +// Driver initialization and reporting: + +#include + +#include "hip_runtime.h" +#include "hcc_detail/hip_hcc.h" +#include "hcc_detail/trace_helper.h" + +// Stack of contexts +thread_local std::stack tls_ctxStack; + + +hipError_t hipInit(unsigned int flags) +{ + HIP_INIT_API(flags); + + hipError_t e = hipSuccess; + + // Flags must be 0 + if (flags != 0) { + e = hipErrorInvalidValue; + } + + return ihipLogStatus(e); +} + + +hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device) +{ + HIP_INIT_API(ctx, flags, device); // FIXME - review if we want to init + hipError_t e = hipSuccess; + + *ctx = new ihipCtx_t(device, g_deviceCnt, flags); + ihipSetTlsDefaultCtx(*ctx); + tls_ctxStack.push(*ctx); + + return ihipLogStatus(e); +} + + +hipError_t hipDeviceGet(hipDevice_t *device, int deviceId) +{ + HIP_INIT_API(device, deviceId); // FIXME - review if we want to init + + *device = ihipGetDevice(deviceId); + + hipError_t e = hipSuccess; + if (*device == NULL) { + e = hipErrorInvalidDevice; + } + + return ihipLogStatus(e); +}; + + +/** + * @return #hipSuccess + */ +//--- +hipError_t hipDriverGetVersion(int *driverVersion) +{ + HIP_INIT_API(driverVersion); + + if (driverVersion) { + *driverVersion = 4; + } + + return ihipLogStatus(hipSuccess); +} + +hipError_t hipCtxDestroy(hipCtx_t ctx) +{ + hipError_t e = hipSuccess; + ihipCtx_t* currentCtx= ihipGetTlsDefaultCtx(); + if(currentCtx == ctx) { + //need to destroy the ctx associated with calling thread + tls_ctxStack.pop(); + } + delete ctx; //As per CUDA docs , attempting to access ctx from those threads which has this ctx as current, will result in the error HIP_ERROR_CONTEXT_IS_DESTROYED. + return ihipLogStatus(e); +} + +hipError_t hipCtxPopCurrent(hipCtx_t* ctx) +{ + hipError_t e = hipSuccess; + ihipCtx_t* tempCtx; + *ctx = ihipGetTlsDefaultCtx(); + if(!tls_ctxStack.empty()) { + tls_ctxStack.pop(); + } + if(!tls_ctxStack.empty()) { + tempCtx= tls_ctxStack.top(); + } + else { + tempCtx = nullptr; + } + + ihipSetTlsDefaultCtx(tempCtx); //TOD0 - Shall check for NULL? + return ihipLogStatus(e); +} + +hipError_t hipCtxPushCurrent(hipCtx_t ctx) +{ + hipError_t e = hipSuccess; + if(ctx != NULL) { //TODO- is this check needed? + ihipSetTlsDefaultCtx(ctx); + tls_ctxStack.push(ctx); + } + else { + e = hipErrorInvalidContext; + } + return ihipLogStatus(e); +} + +hipError_t hipCtxGetCurrent(hipCtx_t* ctx) +{ + hipError_t e = hipSuccess; + + *ctx = ihipGetTlsDefaultCtx(); + if(*ctx == nullptr) { + *ctx = NULL; //TODO - is it required? Can return nullptr? + } + return ihipLogStatus(e); +} + +hipError_t hipCtxSetCurrent(hipCtx_t ctx) +{ + hipError_t e = hipSuccess; + if(ctx == NULL) { + tls_ctxStack.pop(); + } + else { + ihipSetTlsDefaultCtx(ctx); + tls_ctxStack.push(ctx); + } + return ihipLogStatus(e); +} + +hipError_t hipCtxGetDevice(hipDevice_t *device) +{ + hipError_t e = hipSuccess; + + ihipCtx_t *ctx = ihipGetTlsDefaultCtx(); + + if(ctx == nullptr) { + e = hipErrorInvalidContext; + } + else { + *device = (ihipDevice_t*)ctx->getDevice(); + } + return ihipLogStatus(e); +} + +hipError_t hipCtxGetApiVersion (hipCtx_t ctx,int *apiVersion) +{ + HIP_INIT_API(apiVersion); + + if (apiVersion) { + *apiVersion = 4; + } + + return ihipLogStatus(hipSuccess); +} + +hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig ) +{ + HIP_INIT_API(cacheConfig); + + *cacheConfig = hipFuncCachePreferNone; + + return ihipLogStatus(hipSuccess); +} + +hipError_t hipCtxSetCacheConfig ( hipFuncCache cacheConfig ) +{ + HIP_INIT_API(cacheConfig); + + // Nop, AMD does not support variable cache configs. + + return ihipLogStatus(hipSuccess); +} + +hipError_t hipCtxSetSharedMemConfig ( hipSharedMemConfig config ) +{ + HIP_INIT_API(config); + + // Nop, AMD does not support variable shared mem configs. + + return ihipLogStatus(hipSuccess); +} + +hipError_t hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig ) +{ + HIP_INIT_API(pConfig); + + *pConfig = hipSharedMemBankSizeFourByte; + + return ihipLogStatus(hipSuccess); +} \ No newline at end of file diff --git a/projects/hip/src/hip_device.cpp b/projects/hip/src/hip_device.cpp index cfc285427c..e3d7fefa91 100644 --- a/projects/hip/src/hip_device.cpp +++ b/projects/hip/src/hip_device.cpp @@ -26,14 +26,25 @@ THE SOFTWARE. //------------------------------------------------------------------------------------------------- //--- /** - * @return #hipSuccess + * @return #hipSuccess, hipErrorInvalidDevice */ -hipError_t hipGetDevice(int *device) +// TODO - does this initialize HIP runtime? +hipError_t hipGetDevice(int *deviceId) { - HIP_INIT_API(device); + HIP_INIT_API(deviceId); - *device = tls_defaultDevice; - return ihipLogStatus(hipSuccess); + hipError_t e = hipSuccess; + + auto ctx = ihipGetTlsDefaultCtx(); + + if (ctx == nullptr) { + e = hipErrorInvalidDevice; // TODO, check error code. + *deviceId = -1; + } else { + *deviceId = ctx->getDevice()->_deviceId; + } + + return ihipLogStatus(e); } @@ -41,6 +52,7 @@ hipError_t hipGetDevice(int *device) /** * @return #hipSuccess, #hipErrorNoDevice */ +// TODO - does this initialize HIP runtime? hipError_t hipGetDeviceCount(int *count) { HIP_INIT_API(count); @@ -130,13 +142,13 @@ hipError_t hipDeviceGetSharedMemConfig ( hipSharedMemConfig * pConfig ) /** * @return #hipSuccess, #hipErrorInvalidDevice */ -hipError_t hipSetDevice(int device) +hipError_t hipSetDevice(int deviceId) { - HIP_INIT_API(device); - if ((device < 0) || (device >= g_deviceCnt)) { + HIP_INIT_API(deviceId); + if ((deviceId < 0) || (deviceId >= g_deviceCnt)) { return ihipLogStatus(hipErrorInvalidDevice); } else { - tls_defaultDevice = device; + ihipSetTlsDefaultCtx(ihipGetPrimaryCtx(deviceId)); return ihipLogStatus(hipSuccess); } } @@ -150,7 +162,7 @@ hipError_t hipDeviceSynchronize(void) { HIP_INIT_API(); - ihipGetTlsDefaultDevice()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish. + ihipGetTlsDefaultCtx()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish. return ihipLogStatus(hipSuccess); } @@ -164,16 +176,16 @@ hipError_t hipDeviceReset(void) { HIP_INIT_API(); - ihipDevice_t *device = ihipGetTlsDefaultDevice(); + auto *ctx = ihipGetTlsDefaultCtx(); // TODO-HCC // This function currently does a user-level cleanup of known resources. // It could benefit from KFD support to perform a more "nuclear" clean that would include any associated kernel resources and page table entries. - if (device) { - // Release device resources (streams and memory): - device->locked_reset(); + if (ctx) { + // Release ctx resources (streams and memory): + ctx->locked_reset(); } return ihipLogStatus(hipSuccess); @@ -188,7 +200,7 @@ hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) hipError_t e = hipSuccess; - ihipDevice_t * hipDevice = ihipGetDevice(device); + auto * hipDevice = ihipGetDevice(device); hipDeviceProp_t *prop = &hipDevice->_props; if (hipDevice) { switch (attr) { @@ -264,7 +276,7 @@ hipError_t hipGetDeviceProperties(hipDeviceProp_t* props, int device) hipError_t e; - ihipDevice_t * hipDevice = ihipGetDevice(device); + auto * hipDevice = ihipGetDevice(device); if (hipDevice) { // copy saved props *props = hipDevice->_props; @@ -283,15 +295,35 @@ hipError_t hipSetDeviceFlags( unsigned int flags) hipError_t e; - ihipDevice_t * hipDevice = ihipGetDevice(tls_defaultDevice); - if(hipDevice){ - hipDevice->_device_flags = hipDevice->_device_flags | flags; + auto * ctx = ihipGetTlsDefaultCtx(); + + // TODO : does this really OR in the flags or replaces previous flags: + // TODO : Review error handling behavior for this function, it often returns ErrorSetOnActiveProcess + if (ctx) { + ctx->_ctxFlags = ctx->_ctxFlags | flags; e = hipSuccess; - }else{ + } else { e = hipErrorInvalidDevice; } + return ihipLogStatus(e); +}; + + + + +hipError_t hipDeviceGetFromId(hipDevice_t *device, int deviceId) +{ + HIP_INIT_API(device, deviceId); + + hipError_t e = hipSuccess; + + *device = ihipGetDevice(deviceId); + + if (device == nullptr) { + e = hipErrorInvalidDevice; + } + + return ihipLogStatus(e); } - - diff --git a/projects/hip/src/hip_error.cpp b/projects/hip/src/hip_error.cpp index 7c723b1aa2..d9c6dd9aa9 100644 --- a/projects/hip/src/hip_error.cpp +++ b/projects/hip/src/hip_error.cpp @@ -40,12 +40,11 @@ hipError_t hipGetLastError() //--- -hipError_t hipPeakAtLastError() +hipError_t hipPeekAtLastError() { HIP_INIT_API(); - - // peak at last error, but don't reset it. + // peek at last error, but don't reset it. return ihipLogStatus(tls_lastHipError); } diff --git a/projects/hip/src/hip_event.cpp b/projects/hip/src/hip_event.cpp index 1514fc5868..ca30c3c62b 100644 --- a/projects/hip/src/hip_event.cpp +++ b/projects/hip/src/hip_event.cpp @@ -39,7 +39,7 @@ hipError_t ihipEventCreate(hipEvent_t* event, unsigned flags) eh->_stream = NULL; eh->_flags = flags; eh->_timestamp = 0; - eh->_copy_seq_id = 0; + eh->_copySeqId = 0; } else { e = hipErrorInvalidValue; } @@ -79,8 +79,8 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) // If stream == NULL, wait on all queues. // TODO-HCC fix this - is this conservative or still uses device timestamps? // TODO-HCC can we use barrier or event marker to implement better solution? - ihipDevice_t *device = ihipGetTlsDefaultDevice(); - device->locked_syncDefaultStream(true); + ihipCtx_t *ctx = ihipGetTlsDefaultCtx(); + ctx->locked_syncDefaultStream(true); eh->_timestamp = hc::get_system_ticks(); eh->_state = hipEventStatusRecorded; @@ -91,7 +91,7 @@ hipError_t hipEventRecord(hipEvent_t event, hipStream_t stream) eh->_timestamp = 0; eh->_marker = stream->_av.create_marker(); - eh->_copy_seq_id = stream->locked_lastCopySeqId(); + eh->_copySeqId = stream->locked_lastCopySeqId(); return ihipLogStatus(hipSuccess); } @@ -130,12 +130,12 @@ hipError_t hipEventSynchronize(hipEvent_t event) // Created but not actually recorded on any device: return ihipLogStatus(hipSuccess); } else if (eh->_stream == NULL) { - ihipDevice_t *device = ihipGetTlsDefaultDevice(); - device->locked_syncDefaultStream(true); + auto *ctx = ihipGetTlsDefaultCtx(); + ctx->locked_syncDefaultStream(true); return ihipLogStatus(hipSuccess); } else { eh->_marker.wait((eh->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked : hc::hcWaitModeActive); - eh->_stream->locked_reclaimSignals(eh->_copy_seq_id); + eh->_stream->locked_reclaimSignals(eh->_copySeqId); return ihipLogStatus(hipSuccess); } diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index c6c8691419..f0d123b64a 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -44,10 +44,13 @@ THE SOFTWARE. #include "hsa_ext_amd.h" #include "hsakmt.h" -// TODO, re-org header order. -extern const char *ihipErrorString(hipError_t hip_error); #include "hcc_detail/trace_helper.h" + + +//================================================================================================= +//Global variables: +//================================================================================================= const int release = 1; #define MEMCPY_D2H_STAGING_VS_PININPLACE_COPY_THRESHOLD 4194304 @@ -76,48 +79,105 @@ int HIP_VISIBLE_DEVICES = 0; /* Contains a comma-separated sequence of GPU ident int HIP_DISABLE_HW_KERNEL_DEP = 0; int HIP_DISABLE_HW_COPY_DEP = 0; -thread_local int tls_defaultDevice = 0; -thread_local hipError_t tls_lastHipError = hipSuccess; -//================================================================================================= -//Forward Declarations: -//================================================================================================= -bool ihipIsValidDevice(unsigned deviceIndex); std::once_flag hip_initialized; -ihipDevice_t *g_devices; + +// Array of pointers to devices. +ihipDevice_t **g_deviceArray; + + bool g_visible_device = false; unsigned g_deviceCnt; std::vector g_hip_visible_devices; hsa_agent_t g_cpu_agent; +// TODO, remove these if possible: +hsa_agent_t gpu_agent_; +hsa_amd_memory_pool_t gpu_pool_; //================================================================================================= -// Implementation: +// 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; + + + + +//================================================================================================= +// Top-level "free" functions: +//================================================================================================= +static inline bool ihipIsValidDevice(unsigned deviceIndex) +{ + // deviceIndex is unsigned so always > 0 + return (deviceIndex < g_deviceCnt); +} + + +ihipDevice_t * ihipGetDevice(int deviceIndex) +{ + if (ihipIsValidDevice(deviceIndex)) { + return g_deviceArray[deviceIndex]; + } else { + return NULL; + } +} + +ihipCtx_t * ihipGetPrimaryCtx(unsigned deviceIndex) +{ + ihipDevice_t *device = ihipGetDevice(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; +} + + //================================================================================================= // ihipSignal_t: //================================================================================================= // //--- -ihipSignal_t::ihipSignal_t() : _sig_id(0) +ihipSignal_t::ihipSignal_t() : _sigId(0) { - if (hsa_signal_create(0/*value*/, 0, NULL, &_hsa_signal) != HSA_STATUS_SUCCESS) { + if (hsa_signal_create(0/*value*/, 0, NULL, &_hsaSignal) != HSA_STATUS_SUCCESS) { throw ihipException(hipErrorRuntimeMemory); } - //tprintf (DB_SIGNAL, " allocated hsa_signal=%lu\n", (_hsa_signal.handle)); + //tprintf (DB_SIGNAL, " allocated hsa_signal=%lu\n", (_hsaSignal.handle)); } //--- ihipSignal_t::~ihipSignal_t() { - tprintf (DB_SIGNAL, " destroy hsa_signal #%lu (#%lu)\n", (_hsa_signal.handle), _sig_id); - if (hsa_signal_destroy(_hsa_signal) != HSA_STATUS_SUCCESS) { + tprintf (DB_SIGNAL, " destroy hsa_signal #%lu (#%lu)\n", (_hsaSignal.handle), _sigId); + if (hsa_signal_destroy(_hsaSignal) != HSA_STATUS_SUCCESS) { throw ihipException(hipErrorRuntimeOther); } }; @@ -128,11 +188,11 @@ ihipSignal_t::~ihipSignal_t() // ihipStream_t: //================================================================================================= //--- -ihipStream_t::ihipStream_t(unsigned device_index, hc::accelerator_view av, unsigned int flags) : +ihipStream_t::ihipStream_t(ihipCtx_t *ctx, hc::accelerator_view av, unsigned int flags) : _id(0), // will be set by add function. _av(av), _flags(flags), - _device_index(device_index) + _ctx(ctx) { tprintf(DB_SYNC, " streamCreate: stream=%p\n", this); }; @@ -144,7 +204,6 @@ ihipStream_t::~ihipStream_t() } - //--- //TODO - this function is dangerous since it does not propertly account //for younger commands which may be depending on the signals we are reclaiming. @@ -162,10 +221,10 @@ void ihipStream_t::locked_reclaimSignals(SIGSEQNUM sigNum) //--- void ihipStream_t::waitCopy(LockedAccessor_StreamCrit_t &crit, ihipSignal_t *signal) { - SIGSEQNUM sigNum = signal->_sig_id; + SIGSEQNUM sigNum = signal->_sigId; tprintf(DB_SYNC, "waitCopy signal:#%lu\n", sigNum); - hsa_signal_wait_acquire(signal->_hsa_signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); + hsa_signal_wait_acquire(signal->_hsaSignal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); tprintf(DB_SIGNAL, "waitCopy reclaim signal #%lu\n", sigNum); @@ -183,8 +242,8 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty if (! assertQueueEmpty) { tprintf (DB_SYNC, "stream %p wait for queue-empty..\n", this); _av.wait(); - } - + } + if (crit->_last_copy_signal) { tprintf (DB_SYNC, "stream %p wait for lastCopy:#%lu...\n", this, lastCopySeqId(crit) ); this->waitCopy(crit, crit->_last_copy_signal); @@ -209,94 +268,27 @@ void ihipStream_t::locked_wait(bool assertQueueEmpty) }; - -// Recompute the peercnt and the packed _peerAgents whenever a peer is added or deleted. -// The packed _peerAgents can efficiently be used on each memory allocation. -template<> -void ihipDeviceCriticalBase_t::recomputePeerAgents() -{ - _peerCnt = 0; - std::for_each (_peers.begin(), _peers.end(), [this](ihipDevice_t* device) { - _peerAgents[_peerCnt++] = device->_hsa_agent; - }); -} +//============================================================================= -template<> -bool ihipDeviceCriticalBase_t::isPeer(const ihipDevice_t *peer) -{ - auto match = std::find(_peers.begin(), _peers.end(), peer); - return (match != std::end(_peers)); -} - - -template<> -bool ihipDeviceCriticalBase_t::addPeer(ihipDevice_t *peer) -{ - auto match = std::find(_peers.begin(), _peers.end(), peer); - if (match == std::end(_peers)) { - // Not already a peer, let's update the list: - _peers.push_back(peer); - recomputePeerAgents(); - return true; - } - - // If we get here - peer was already on list, silently ignore. - return false; -} - - -template<> -bool ihipDeviceCriticalBase_t::removePeer(ihipDevice_t *peer) -{ - auto match = std::find(_peers.begin(), _peers.end(), peer); - if (match != std::end(_peers)) { - // Found a valid peer, let's remove it. - _peers.remove(peer); - recomputePeerAgents(); - return true; - } else { - return false; - } -} - - -template<> -void ihipDeviceCriticalBase_t::resetPeers(ihipDevice_t *thisDevice) -{ - _peers.clear(); - _peerCnt = 0; - addPeer(thisDevice); // peer-list always contains self agent. -} - - -template<> -void ihipDeviceCriticalBase_t::addStream(ihipStream_t *stream) -{ - _streams.push_back(stream); - stream->_id = incStreamId(); -} //------------------------------------------------------------------------------------------------- + //--- -//Flavor that takes device index. -ihipDevice_t * getDevice(unsigned deviceIndex) +const ihipDevice_t * ihipStream_t::getDevice() const { - if (ihipIsValidDevice(deviceIndex)) { - return &g_devices[deviceIndex]; - } else { - return NULL; - } + return _ctx->getDevice(); }; -//--- -ihipDevice_t * ihipStream_t::getDevice() const + +ihipCtx_t * ihipStream_t::getCtx() const { - return ::getDevice(_device_index); + return _ctx; }; + #define HIP_NUM_SIGNALS_PER_STREAM 32 @@ -322,12 +314,12 @@ ihipSignal_t *ihipStream_t::allocSignal(LockedAccessor_StreamCrit_t &crit) crit->_signalCursor = 0; } - if (crit->_signalPool[thisCursor]._sig_id < crit->_oldest_live_sig_id) { - SIGSEQNUM oldSigId = crit->_signalPool[thisCursor]._sig_id; + if (crit->_signalPool[thisCursor]._sigId < crit->_oldest_live_sig_id) { + SIGSEQNUM oldSigId = crit->_signalPool[thisCursor]._sigId; crit->_signalPool[thisCursor]._index = thisCursor; - crit->_signalPool[thisCursor]._sig_id = ++crit->_stream_sig_id; // allocate it. + crit->_signalPool[thisCursor]._sigId = ++crit->_streamSigId; // allocate it. tprintf(DB_SIGNAL, "allocatSignal #%lu at pos:%i (old sigId:%lu < oldest_live:%lu)\n", - crit->_signalPool[thisCursor]._sig_id, + crit->_signalPool[thisCursor]._sigId, thisCursor, oldSigId, crit->_oldest_live_sig_id); @@ -372,9 +364,9 @@ void ihipStream_t::enqueueBarrier(hsa_queue_t* queue, ihipSignal_t *depSignal, i //header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE; barrier->header = header; - barrier->dep_signal[0].handle = depSignal ? depSignal->_hsa_signal.handle: 0; + barrier->dep_signal[0].handle = depSignal ? depSignal->_hsaSignal.handle: 0; - barrier->completion_signal.handle = completionSignal ? completionSignal->_hsa_signal.handle : 0; + barrier->completion_signal.handle = completionSignal ? completionSignal->_hsaSignal.handle : 0; // TODO - check queue overflow, return error: // Increment write index and ring doorbell to dispatch the kernel @@ -411,7 +403,7 @@ bool ihipStream_t::lockopen_preKernelCommand() if (HIP_DISABLE_HW_KERNEL_DEP == 0) { this->enqueueBarrier(q, crit->_last_copy_signal, NULL); tprintf (DB_SYNC, "stream %p switch %s to %s (barrier pkt inserted with wait on #%lu)\n", - this, ihipCommandName[crit->_last_command_type], ihipCommandName[ihipCommandKernel], crit->_last_copy_signal->_sig_id) + this, ihipCommandName[crit->_last_command_type], ihipCommandName[ihipCommandKernel], crit->_last_copy_signal->_sigId) } else if (HIP_DISABLE_HW_KERNEL_DEP>0) { tprintf (DB_SYNC, "stream %p switch %s to %s (HOST wait for previous...)\n", @@ -459,14 +451,14 @@ int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t this, ihipCommandName[crit->_last_command_type], ihipCommandName[copyType]); needSync = 1; ihipSignal_t *depSignal = allocSignal(crit); - hsa_signal_store_relaxed(depSignal->_hsa_signal,1); + hsa_signal_store_relaxed(depSignal->_hsaSignal,1); this->enqueueBarrier(static_cast(_av.get_hsa_queue()), NULL, depSignal); - *waitSignal = depSignal->_hsa_signal; + *waitSignal = depSignal->_hsaSignal; } else if (crit->_last_copy_signal) { needSync = 1; tprintf (DB_SYNC, "stream %p switch %s to %s (async copy dep on other copy #%lu)\n", - this, ihipCommandName[crit->_last_command_type], ihipCommandName[copyType], crit->_last_copy_signal->_sig_id); - *waitSignal = crit->_last_copy_signal->_hsa_signal; + this, ihipCommandName[crit->_last_command_type], ihipCommandName[copyType], crit->_last_copy_signal->_sigId); + *waitSignal = crit->_last_copy_signal->_hsaSignal; } if (HIP_DISABLE_HW_COPY_DEP && needSync) { @@ -491,115 +483,115 @@ int ihipStream_t::preCopyCommand(LockedAccessor_StreamCrit_t &crit, ihipSignal_t +//============================================================================= +// Recompute the peercnt and the packed _peerAgents whenever a peer is added or deleted. +// The packed _peerAgents can efficiently be used on each memory allocation. +template<> +void ihipCtxCriticalBase_t::recomputePeerAgents() +{ + _peerCnt = 0; + std::for_each (_peers.begin(), _peers.end(), [this](ihipCtx_t* ctx) { + _peerAgents[_peerCnt++] = ctx->getDevice()->_hsaAgent; + }); +} + + +template<> +bool ihipCtxCriticalBase_t::isPeer(const ihipCtx_t *peer) +{ + auto match = std::find(_peers.begin(), _peers.end(), peer); + return (match != std::end(_peers)); +} + + +template<> +bool ihipCtxCriticalBase_t::addPeer(ihipCtx_t *peer) +{ + auto match = std::find(_peers.begin(), _peers.end(), peer); + if (match == std::end(_peers)) { + // Not already a peer, let's update the list: + _peers.push_back(peer); + recomputePeerAgents(); + return true; + } + + // If we get here - peer was already on list, silently ignore. + return false; +} + + +template<> +bool ihipCtxCriticalBase_t::removePeer(ihipCtx_t *peer) +{ + auto match = std::find(_peers.begin(), _peers.end(), peer); + if (match != std::end(_peers)) { + // Found a valid peer, let's remove it. + _peers.remove(peer); + recomputePeerAgents(); + return true; + } else { + return false; + } +} + + +template<> +void ihipCtxCriticalBase_t::resetPeers(ihipCtx_t *thisDevice) +{ + _peers.clear(); + _peerCnt = 0; + addPeer(thisDevice); // peer-list always contains self agent. +} + + +template<> +void ihipCtxCriticalBase_t::addStream(ihipStream_t *stream) +{ + stream->_id = _streams.size(); + _streams.push_back(stream); +} +//============================================================================= //================================================================================================= -// -//Reset the device - this is called from hipDeviceReset. -//Device may be reset multiple times, and may be reset after init. -void ihipDevice_t::locked_reset() +// ihipDevice_t +//================================================================================================= +ihipDevice_t::ihipDevice_t(unsigned deviceId, unsigned deviceCnt, hc::accelerator &acc) : + _deviceId(deviceId), + _acc(acc) { - // Obtain mutex access to the device critical data, release by destructor - LockedAccessor_DeviceCrit_t crit(_criticalData); - - - //--- - //Wait for pending activity to complete? TODO - check if this is required behavior: - tprintf(DB_SYNC, "locked_reset waiting for activity to complete.\n"); - - // Reset and remove streams: - // Delete all created streams including the default one. - for (auto streamI=crit->const_streams().begin(); streamI!=crit->const_streams().end(); streamI++) { - ihipStream_t *stream = *streamI; - (*streamI)->locked_wait(); - tprintf(DB_SYNC, " delete stream=%p\n", stream); - - delete stream; - } - // Clear the list. - crit->streams().clear(); - - - // Create a fresh default stream and add it: - _default_stream = new ihipStream_t(_device_index, _acc.get_default_view(), hipStreamDefault); - crit->addStream(_default_stream); - - - // This resest peer list to just me: - crit->resetPeers(this); - - // Reset and release all memory stored in the tracker: - // Reset will remove peer mapping so don't need to do this explicitly. - am_memtracker_reset(_acc); - -}; - - -//--- -void ihipDevice_t::init(unsigned device_index, unsigned deviceCnt, hc::accelerator &acc, unsigned flags) -{ - _device_index = device_index; - _device_flags = flags; - _acc = acc; - hsa_agent_t *agent = static_cast (acc.get_hsa_agent()); if (agent) { - int err = hsa_agent_get_info(*agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT, &_compute_units); + int err = hsa_agent_get_info(*agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT, &_computeUnits); if (err != HSA_STATUS_SUCCESS) { - _compute_units = 1; + _computeUnits = 1; } - _hsa_agent = *agent; + _hsaAgent = *agent; } else { - _hsa_agent.handle = static_cast (-1); + _hsaAgent.handle = static_cast (-1); } - getProperties(&_props); - - _criticalData.init(deviceCnt); - - locked_reset(); - - - tprintf(DB_SYNC, "created device with default_stream=%p\n", _default_stream); - - hsa_region_t *pinnedHostRegion; - pinnedHostRegion = static_cast(_acc.get_hsa_am_system_region()); - _staging_buffer[0] = new StagingBuffer(_hsa_agent, *pinnedHostRegion, HIP_STAGING_SIZE*1024, HIP_STAGING_BUFFERS); - _staging_buffer[1] = new StagingBuffer(_hsa_agent, *pinnedHostRegion, HIP_STAGING_SIZE*1024, HIP_STAGING_BUFFERS); - -}; + initProperties(&_props); + _stagingBuffer[0] = new UnpinnedCopyEngine(_hsaAgent,g_cpu_agent, HIP_STAGING_SIZE*1024, HIP_STAGING_BUFFERS,HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING,HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE,HIP_D2H_MEM_TRANSFER_THRESHOLD); + _stagingBuffer[1] = new UnpinnedCopyEngine(_hsaAgent,g_cpu_agent, HIP_STAGING_SIZE*1024, HIP_STAGING_BUFFERS,HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING,HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE,HIP_D2H_MEM_TRANSFER_THRESHOLD); + _primaryCtx = new ihipCtx_t(this, deviceCnt, hipDeviceMapHost); +} ihipDevice_t::~ihipDevice_t() { - if (_default_stream) { - delete _default_stream; - _default_stream = NULL; - } - for (int i=0; i<2; i++) { - if (_staging_buffer[i]) { - delete _staging_buffer[i]; - _staging_buffer[i] = NULL; + if (_stagingBuffer[i]) { + delete _stagingBuffer[i]; + _stagingBuffer[i] = NULL; } } } -//---- - - -//================================================================================================= -// Utility functions, these are not part of the public HIP API -//================================================================================================= - -//================================================================================================= - -#define DeviceErrorCheck(x) if (x != HSA_STATUS_SUCCESS) { return hipErrorInvalidDevice; } - #define ErrorCheck(x) error_check(x, __LINE__, __FILE__) void error_check(hsa_status_t hsa_error_code, int line_num, std::string str) { @@ -608,13 +600,28 @@ void error_check(hsa_status_t hsa_error_code, int line_num, std::string str) { } } -// CPU agent used for verification -hsa_agent_t cpu_agent_; -hsa_agent_t gpu_agent_; -int gpu_region_count; -// System region -hsa_amd_memory_pool_t sys_region_; -hsa_amd_memory_pool_t gpu_region_; + + +//--- +// Helper for initProperties +// Determines if the given agent is of type HSA_DEVICE_TYPE_GPU and counts it. +static hsa_status_t countGpuAgents(hsa_agent_t agent, void *data) { + if (data == NULL) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } + hsa_device_type_t device_type; + hsa_status_t status = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &device_type); + if (status != HSA_STATUS_SUCCESS) { + return status; + } + if (device_type == HSA_DEVICE_TYPE_GPU) { + (*static_cast(data))++; + } + return HSA_STATUS_SUCCESS; +} + + + hsa_status_t FindGpuDevice(hsa_agent_t agent, void* data) { if (data == NULL) { @@ -636,27 +643,7 @@ hsa_status_t FindGpuDevice(hsa_agent_t agent, void* data) { return HSA_STATUS_SUCCESS; } -hsa_status_t FindCpuDevice(hsa_agent_t agent, void* data) { - if (data == NULL) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } - - hsa_device_type_t hsa_device_type; - hsa_status_t hsa_error_code = - hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &hsa_device_type); - if (hsa_error_code != HSA_STATUS_SUCCESS) { - return hsa_error_code; - } - - if (hsa_device_type == HSA_DEVICE_TYPE_CPU) { - *((hsa_agent_t*)data) = agent; - return HSA_STATUS_INFO_BREAK; - } - - return HSA_STATUS_SUCCESS; -} - -hsa_status_t GetDeviceRegion(hsa_amd_memory_pool_t region, void* data) { +hsa_status_t GetDevicePool(hsa_amd_memory_pool_t pool, void* data) { if (NULL == data) { return HSA_STATUS_ERROR_INVALID_ARGUMENT; } @@ -665,50 +652,21 @@ hsa_status_t GetDeviceRegion(hsa_amd_memory_pool_t region, void* data) { hsa_amd_segment_t segment; uint32_t flag; - err = hsa_amd_memory_pool_get_info(region, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment); + err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment); ErrorCheck(err); if (HSA_AMD_SEGMENT_GLOBAL != segment) return HSA_STATUS_SUCCESS; - err = hsa_amd_memory_pool_get_info(region, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag); + err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag); ErrorCheck(err); - *((hsa_amd_memory_pool_t*)data) = region; + *((hsa_amd_memory_pool_t*)data) = pool; return HSA_STATUS_SUCCESS; } -hsa_status_t FindGlobalRegion(hsa_amd_memory_pool_t region, void* data) { - if (NULL == data) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } - - hsa_status_t err; - hsa_amd_segment_t segment; - uint32_t flag; - err = hsa_amd_memory_pool_get_info(region, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment); - ErrorCheck(err); - - err = hsa_amd_memory_pool_get_info(region, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag); - ErrorCheck(err); - if ((HSA_AMD_SEGMENT_GLOBAL == segment) && - (flag & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_FINE_GRAINED)) { - *((hsa_amd_memory_pool_t*)data) = region; - } - return HSA_STATUS_SUCCESS; -} - -void FindDeviceRegion() +void FindDevicePool() { hsa_status_t err = hsa_iterate_agents(FindGpuDevice, &gpu_agent_); ErrorCheck(err); - err = hsa_amd_agent_iterate_memory_pools(gpu_agent_, GetDeviceRegion, &gpu_region_); - ErrorCheck(err); -} - -void FindSystemRegion() -{ - hsa_status_t err = hsa_iterate_agents(FindCpuDevice, &cpu_agent_); - ErrorCheck(err); - - err = hsa_amd_agent_iterate_memory_pools(cpu_agent_, FindGlobalRegion, &sys_region_); + err = hsa_amd_agent_iterate_memory_pools(gpu_agent_, GetDevicePool, &gpu_pool_); ErrorCheck(err); } @@ -743,24 +701,30 @@ hsa_status_t get_region_info(hsa_region_t region, void* data) return HSA_STATUS_SUCCESS; } + // Determines if the given agent is of type HSA_DEVICE_TYPE_GPU and counts it. -static hsa_status_t countGpuAgents(hsa_agent_t agent, void *data) { - if (data == NULL) { - return HSA_STATUS_ERROR_INVALID_ARGUMENT; - } +static hsa_status_t findCpuAgent(hsa_agent_t agent, void *data) +{ hsa_device_type_t device_type; hsa_status_t status = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &device_type); if (status != HSA_STATUS_SUCCESS) { return status; } - if (device_type == HSA_DEVICE_TYPE_GPU) { - (*static_cast(data))++; + if (device_type == HSA_DEVICE_TYPE_CPU) { + (*static_cast(data)) = agent; + return HSA_STATUS_INFO_BREAK; } + return HSA_STATUS_SUCCESS; } -// Internal version, -hipError_t ihipDevice_t::getProperties(hipDeviceProp_t* prop) + +#define DeviceErrorCheck(x) if (x != HSA_STATUS_SUCCESS) { return hipErrorInvalidDevice; } + +//--- +// Initialize properties for the device. +// Call this once when the ihipDevice_t is created: +hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) { hipError_t e = hipSuccess; hsa_status_t err; @@ -772,7 +736,7 @@ hipError_t ihipDevice_t::getProperties(hipDeviceProp_t* prop) prop-> maxThreadsPerMultiProcessor = 0; prop->regsPerBlock = 0; - if (_hsa_agent.handle == -1) { + if (_hsaAgent.handle == -1) { return hipErrorInvalidDevice; } @@ -786,39 +750,39 @@ hipError_t ihipDevice_t::getProperties(hipDeviceProp_t* prop) prop->isMultiGpuBoard = 0 ? gpuAgentsCount < 2 : 1; // Get agent name - err = hsa_agent_get_info(_hsa_agent, HSA_AGENT_INFO_NAME, &(prop->name)); + err = hsa_agent_get_info(_hsaAgent, HSA_AGENT_INFO_NAME, &(prop->name)); DeviceErrorCheck(err); // Get agent node uint32_t node; - err = hsa_agent_get_info(_hsa_agent, HSA_AGENT_INFO_NODE, &node); + err = hsa_agent_get_info(_hsaAgent, HSA_AGENT_INFO_NODE, &node); DeviceErrorCheck(err); // Get wavefront size - err = hsa_agent_get_info(_hsa_agent, HSA_AGENT_INFO_WAVEFRONT_SIZE,&prop->warpSize); + err = hsa_agent_get_info(_hsaAgent, HSA_AGENT_INFO_WAVEFRONT_SIZE,&prop->warpSize); DeviceErrorCheck(err); // Get max total number of work-items in a workgroup - err = hsa_agent_get_info(_hsa_agent, HSA_AGENT_INFO_WORKGROUP_MAX_SIZE, &prop->maxThreadsPerBlock ); + err = hsa_agent_get_info(_hsaAgent, HSA_AGENT_INFO_WORKGROUP_MAX_SIZE, &prop->maxThreadsPerBlock ); DeviceErrorCheck(err); // Get max number of work-items of each dimension of a work-group uint16_t work_group_max_dim[3]; - err = hsa_agent_get_info(_hsa_agent, HSA_AGENT_INFO_WORKGROUP_MAX_DIM, work_group_max_dim); + err = hsa_agent_get_info(_hsaAgent, HSA_AGENT_INFO_WORKGROUP_MAX_DIM, work_group_max_dim); DeviceErrorCheck(err); for( int i =0; i< 3 ; i++) { prop->maxThreadsDim[i]= work_group_max_dim[i]; } hsa_dim3_t grid_max_dim; - err = hsa_agent_get_info(_hsa_agent, HSA_AGENT_INFO_GRID_MAX_DIM, &grid_max_dim); + err = hsa_agent_get_info(_hsaAgent, HSA_AGENT_INFO_GRID_MAX_DIM, &grid_max_dim); DeviceErrorCheck(err); prop->maxGridSize[0]= (int) ((grid_max_dim.x == UINT32_MAX) ? (INT32_MAX) : grid_max_dim.x); prop->maxGridSize[1]= (int) ((grid_max_dim.y == UINT32_MAX) ? (INT32_MAX) : grid_max_dim.y); prop->maxGridSize[2]= (int) ((grid_max_dim.z == UINT32_MAX) ? (INT32_MAX) : grid_max_dim.z); // Get Max clock frequency - err = hsa_agent_get_info(_hsa_agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_MAX_CLOCK_FREQUENCY, &prop->clockRate); + err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_MAX_CLOCK_FREQUENCY, &prop->clockRate); prop->clockRate *= 1000.0; // convert Mhz to Khz. DeviceErrorCheck(err); @@ -830,7 +794,7 @@ hipError_t ihipDevice_t::getProperties(hipDeviceProp_t* prop) // Get Agent BDFID (bus/device/function ID) uint16_t bdf_id = 1; - err = hsa_agent_get_info(_hsa_agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_BDFID, &bdf_id); + err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_BDFID, &bdf_id); DeviceErrorCheck(err); // BDFID is 16bit uint: [8bit - BusID | 5bit - Device ID | 3bit - Function/DomainID] @@ -845,26 +809,24 @@ hipError_t ihipDevice_t::getProperties(hipDeviceProp_t* prop) prop->minor = 0; // Get number of Compute Unit - err = hsa_agent_get_info(_hsa_agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT, &(prop->multiProcessorCount)); + err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT, &(prop->multiProcessorCount)); DeviceErrorCheck(err); // TODO-hsart - this appears to return 0? uint32_t cache_size[4]; - err = hsa_agent_get_info(_hsa_agent, HSA_AGENT_INFO_CACHE_SIZE, cache_size); + err = hsa_agent_get_info(_hsaAgent, HSA_AGENT_INFO_CACHE_SIZE, cache_size); DeviceErrorCheck(err); prop->l2CacheSize = cache_size[1]; /* Computemode for HSA Devices is always : cudaComputeModeDefault */ prop->computeMode = 0; - FindSystemRegion(); - FindDeviceRegion(); - int access=checkAccess(cpu_agent_, gpu_region_); - if(0!= access){ - isLargeBar= 1; - } - else{ - isLargeBar=0; + FindDevicePool(); + int access=checkAccess(g_cpu_agent, gpu_pool_); + if (0!= access){ + _isLargeBar= 1; + } else { + _isLargeBar=0; } @@ -883,7 +845,7 @@ hipError_t ihipDevice_t::getProperties(hipDeviceProp_t* prop) // Get memory properties - err = hsa_agent_iterate_regions(_hsa_agent, get_region_info, prop); + err = hsa_agent_iterate_regions(_hsaAgent, get_region_info, prop); DeviceErrorCheck(err); // Get the size of the region we are using for Accelerator Memory allocations: @@ -925,22 +887,110 @@ hipError_t ihipDevice_t::getProperties(hipDeviceProp_t* prop) prop->arch.has3dGrid = 1; prop->arch.hasDynamicParallelism = 0; - prop->concurrentKernels = 1; // All ROCR hardware supports executing multiple kernels concurrently + prop->concurrentKernels = 1; // All ROCm hardware supports executing multiple kernels concurrently + + prop->canMapHostMemory = 1; // All ROCm devices can map host memory +#if 0 + // TODO - code broken below since it always returns 1. + // Are the flags part of the context or part of the device? if ( _device_flags | hipDeviceMapHost) { prop->canMapHostMemory = 1; } else { prop->canMapHostMemory = 0; } +#endif return e; } +//================================================================================================= +// ihipCtx_t +//================================================================================================= +ihipCtx_t::ihipCtx_t(ihipDevice_t *device, unsigned deviceCnt, unsigned flags) : + _ctxFlags(flags), + _device(device), + _criticalData(deviceCnt) +{ + locked_reset(); + + tprintf(DB_SYNC, "created ctx with defaultStream=%p\n", _defaultStream); +}; + + + + +ihipCtx_t::~ihipCtx_t() +{ + if (_defaultStream) { + delete _defaultStream; + _defaultStream = NULL; + } +} +//Reset the device - this is called from hipDeviceReset. +//Device may be reset multiple times, and may be reset after init. +void ihipCtx_t::locked_reset() +{ + // Obtain mutex access to the device critical data, release by destructor + LockedAccessor_CtxCrit_t crit(_criticalData); + + + //--- + //Wait for pending activity to complete? TODO - check if this is required behavior: + tprintf(DB_SYNC, "locked_reset waiting for activity to complete.\n"); + + // Reset and remove streams: + // Delete all created streams including the default one. + for (auto streamI=crit->const_streams().begin(); streamI!=crit->const_streams().end(); streamI++) { + ihipStream_t *stream = *streamI; + (*streamI)->locked_wait(); + tprintf(DB_SYNC, " delete stream=%p\n", stream); + + delete stream; + } + // Clear the list. + crit->streams().clear(); + + + // Create a fresh default stream and add it: + _defaultStream = new ihipStream_t(this, getDevice()->_acc.get_default_view(), hipStreamDefault); + crit->addStream(_defaultStream); + + + // Reset peer list to just me: + crit->resetPeers(this); + + // Reset and release all memory stored in the tracker: + // Reset will remove peer mapping so don't need to do this explicitly. + // FIXME - This is clearly a non-const action! Is this a context reset or a device reset - maybe should reference count? + ihipDevice_t *device = getWriteableDevice(); + am_memtracker_reset(device->_acc); + +}; + + + +//---- + + + + +//================================================================================================= +// Utility functions, these are not part of the public HIP API +//================================================================================================= + +//================================================================================================= + + + + + + // Implement "default" stream syncronization // This waits for all other streams to drain before continuing. // If waitOnSelf is set, this additionally waits for the default stream to empty. -void ihipDevice_t::locked_syncDefaultStream(bool waitOnSelf) +void ihipCtx_t::locked_syncDefaultStream(bool waitOnSelf) { - LockedAccessor_DeviceCrit_t crit(_criticalData); + LockedAccessor_CtxCrit_t crit(_criticalData); tprintf(DB_SYNC, "syncDefaultStream\n"); @@ -951,7 +1001,7 @@ void ihipDevice_t::locked_syncDefaultStream(bool waitOnSelf) // And - don't wait for the NULL stream if (!(stream->_flags & hipStreamNonBlocking)) { - if (waitOnSelf || (stream != _default_stream)) { + if (waitOnSelf || (stream != _defaultStream)) { // TODO-hcc - use blocking or active wait here? // TODO-sync - cudaDeviceBlockingSync stream->locked_wait(); @@ -961,17 +1011,17 @@ void ihipDevice_t::locked_syncDefaultStream(bool waitOnSelf) } //--- -void ihipDevice_t::locked_addStream(ihipStream_t *s) +void ihipCtx_t::locked_addStream(ihipStream_t *s) { - LockedAccessor_DeviceCrit_t crit(_criticalData); + LockedAccessor_CtxCrit_t crit(_criticalData); crit->addStream(s); } //--- -void ihipDevice_t::locked_removeStream(ihipStream_t *s) +void ihipCtx_t::locked_removeStream(ihipStream_t *s) { - LockedAccessor_DeviceCrit_t crit(_criticalData); + LockedAccessor_CtxCrit_t crit(_criticalData); crit->streams().remove(s); } @@ -979,9 +1029,9 @@ void ihipDevice_t::locked_removeStream(ihipStream_t *s) //--- //Heavyweight synchronization that waits on all streams, ignoring hipStreamNonBlocking flag. -void ihipDevice_t::locked_waitAllStreams() +void ihipCtx_t::locked_waitAllStreams() { - LockedAccessor_DeviceCrit_t crit(_criticalData); + LockedAccessor_CtxCrit_t crit(_criticalData); tprintf(DB_SYNC, "waitAllStream\n"); for (auto streamI=crit->const_streams().begin(); streamI!=crit->const_streams().end(); streamI++) { @@ -1056,21 +1106,6 @@ void ihipReadEnv_I(int *var_ptr, const char *var_name1, const char *var_name2, c #endif -// Determines if the given agent is of type HSA_DEVICE_TYPE_GPU and counts it. -static hsa_status_t findCpuAgent(hsa_agent_t agent, void *data) -{ - hsa_device_type_t device_type; - hsa_status_t status = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &device_type); - if (status != HSA_STATUS_SUCCESS) { - return status; - } - if (device_type == HSA_DEVICE_TYPE_CPU) { - (*static_cast(data)) = agent; - return HSA_STATUS_INFO_BREAK; - } - - return HSA_STATUS_SUCCESS; -} //--- @@ -1166,7 +1201,13 @@ void ihipInit() } } - g_devices = new ihipDevice_t[deviceCnt]; + hsa_status_t err = hsa_iterate_agents(findCpuAgent, &g_cpu_agent); + if (err != HSA_STATUS_INFO_BREAK) { + // didn't find a CPU. + throw ihipException(hipErrorRuntimeOther); + } + + g_deviceArray = new ihipDevice_t* [deviceCnt]; g_deviceCnt = 0; for (int i=0; i"); } -bool ihipIsValidDevice(unsigned deviceIndex) -{ - // deviceIndex is unsigned so always > 0 - return (deviceIndex < g_deviceCnt); -} - -//--- -ihipDevice_t *ihipGetTlsDefaultDevice() -{ - // If this is invalid, the TLS state is corrupt. - // This can fire if called before devices are initialized. - // TODO - consider replacing assert with error code - assert (ihipIsValidDevice(tls_defaultDevice)); - - return &g_devices[tls_defaultDevice]; -} -//--- -ihipDevice_t *ihipGetDevice(int deviceId) -{ - if ((deviceId >= 0) && (deviceId < g_deviceCnt)) { - return &g_devices[deviceId]; - } else { - return NULL; - } - -} //--- // Get the stream to use for a command submission. @@ -1235,17 +1243,17 @@ ihipDevice_t *ihipGetDevice(int deviceId) hipStream_t ihipSyncAndResolveStream(hipStream_t stream) { if (stream == hipStreamNull ) { - ihipDevice_t *device = ihipGetTlsDefaultDevice(); + ihipCtx_t *device = ihipGetTlsDefaultCtx(); #ifndef HIP_API_PER_THREAD_DEFAULT_STREAM device->locked_syncDefaultStream(false); #endif - return device->_default_stream; + return device->_defaultStream; } else { // Have to wait for legacy default stream to be empty: if (!(stream->_flags & hipStreamNonBlocking)) { tprintf(DB_SYNC, "stream %p wait default stream\n", stream); - stream->getDevice()->_default_stream->locked_wait(); + stream->getCtx()->_defaultStream->locked_wait(); } return stream; @@ -1260,7 +1268,7 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_ { HIP_INIT_API(stream, grid, block, lp); stream = ihipSyncAndResolveStream(stream); -#if USE_GRID_LAUNCH_20 +#if USE_GRID_LAUNCH_20 lp->grid_dim.x = grid.x; lp->grid_dim.y = grid.y; lp->grid_dim.z = grid.z; @@ -1289,7 +1297,7 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, gri { HIP_INIT_API(stream, grid, block, lp); stream = ihipSyncAndResolveStream(stream); -#if USE_GRID_LAUNCH_20 +#if USE_GRID_LAUNCH_20 lp->grid_dim.x = grid; lp->grid_dim.y = 1; lp->grid_dim.z = 1; @@ -1319,7 +1327,7 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, gri { HIP_INIT_API(stream, grid, block, lp); stream = ihipSyncAndResolveStream(stream); -#if USE_GRID_LAUNCH_20 +#if USE_GRID_LAUNCH_20 lp->grid_dim.x = grid.x; lp->grid_dim.y = grid.y; lp->grid_dim.z = grid.z; @@ -1349,7 +1357,7 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, g { HIP_INIT_API(stream, grid, block, lp); stream = ihipSyncAndResolveStream(stream); -#if USE_GRID_LAUNCH_20 +#if USE_GRID_LAUNCH_20 lp->grid_dim.x = grid; lp->grid_dim.y = 1; lp->grid_dim.z = 1; @@ -1479,12 +1487,12 @@ unsigned ihipStream_t::resolveMemcpyDirection(bool srcTracked, bool dstTracked, // Setup the copyCommandType and the copy agents (for hsa_amd_memory_async_copy) -// srcPhysAcc is the physical location of the src data. For many copies this is the +// srcPhysAcc is the physical location of the src data. For many copies this is the void ihipStream_t::setAsyncCopyAgents(unsigned kind, ihipCommand_t *commandType, hsa_agent_t *srcAgent, hsa_agent_t *dstAgent) { // current* represents the device associated with the specified stream. - ihipDevice_t *streamDevice = this->getDevice(); - hsa_agent_t streamAgent = streamDevice->_hsa_agent; + const ihipDevice_t *streamDevice = this->getDevice(); + hsa_agent_t streamAgent = streamDevice->_hsaAgent; // ROCR runtime logic is : // - If both src and dst are cpu agent, launch thread and memcpy. We want to avoid this. @@ -1503,7 +1511,9 @@ void ihipStream_t::setAsyncCopyAgents(unsigned kind, ihipCommand_t *commandType, void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const void* src, size_t sizeBytes, unsigned kind) { - ihipDevice_t *device = this->getDevice(); + ihipCtx_t *ctx = this->getCtx(); + const ihipDevice_t *device = ctx->getDevice(); + if (device == NULL) { throw ihipException(hipErrorInvalidDevice); } @@ -1527,9 +1537,11 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const bool copyEngineCanSeeSrcAndDest = false; if (kind == hipMemcpyDeviceToDevice) { #if USE_PEER_TO_PEER>=2 - // TODO - consider refactor. Do we need to support simul access of enable/disable peers with access? - LockedAccessor_DeviceCrit_t dcrit(device->criticalData()); - if (dcrit->isPeer(::getDevice(dstPtrInfo._appId)) && (dcrit->isPeer(::getDevice(srcPtrInfo._appId)))) { + // Lock to prevent another thread from modifying peer list while we are trying to look at it. + LockedAccessor_CtxCrit_t dcrit(ctx->criticalData()); + // FIXME - this assumes peer access only from primary context. + // Would need to change the tracker to store a void * parameter that we could map to the ctx where the pointer is allocated. + if (dcrit->isPeer(ihipGetPrimaryCtx(dstPtrInfo._appId)) && (dcrit->isPeer(ihipGetPrimaryCtx(srcPtrInfo._appId)))) { copyEngineCanSeeSrcAndDest = true; } #endif @@ -1542,28 +1554,17 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const tprintf(DB_COPY1, "D2H && !dstTracked: staged copy H2D dst=%p src=%p sz=%zu\n", dst, src, sizeBytes); if(HIP_OPTIMAL_MEM_TRANSFER) { - if((device->isLargeBar)&&(sizeBytes < HIP_H2D_MEM_TRANSFER_THRESHOLD_DIRECT_OR_STAGING)){ - memcpy(dst,src,sizeBytes); - std::atomic_thread_fence(std::memory_order_release); - } - else{ - if(sizeBytes > HIP_H2D_MEM_TRANSFER_THRESHOLD_STAGING_OR_PININPLACE){ - //if (HIP_PININPLACE) { - device->_staging_buffer[0]->CopyHostToDevicePinInPlace(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); - } else { - device->_staging_buffer[0]->CopyHostToDevice(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); - } - // The copy waits for inputs and then completes before returning so can reset queue to empty: - this->wait(crit, true); - } + device->_stagingBuffer[0]->CopyHostToDevice(1,device->_isLargeBar,dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); } else { if (HIP_PININPLACE) { - device->_staging_buffer[0]->CopyHostToDevicePinInPlace(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); + device->_stagingBuffer[0]->CopyHostToDevicePinInPlace(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); } else { - device->_staging_buffer[0]->CopyHostToDevice(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); + device->_stagingBuffer[0]->CopyHostToDevice(0,0,dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); } } + // The copy waits for inputs and then completes before returning so can reset queue to empty: + this->wait(crit, true); } else { // TODO - remove, slow path. @@ -1580,7 +1581,7 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const hsa_agent_t srcAgent = *(static_cast(srcPtrInfo._acc.get_hsa_agent())); ihipSignal_t *ihipSignal = allocSignal(crit); - hsa_signal_t copyCompleteSignal = ihipSignal->_hsa_signal; + hsa_signal_t copyCompleteSignal = ihipSignal->_hsaSignal; hsa_signal_store_relaxed(copyCompleteSignal, 1); void *devPtrSrc = srcPtrInfo._devicePointer; @@ -1603,15 +1604,12 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const //printf ("staged-copy- read dep signals\n"); if(HIP_OPTIMAL_MEM_TRANSFER) { - if(sizeBytes> HIP_D2H_MEM_TRANSFER_THRESHOLD){ - device->_staging_buffer[1]->CopyDeviceToHostPinInPlace(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); - }else { - //printf ("staged-copy- read dep signals\n"); - device->_staging_buffer[1]->CopyDeviceToHost(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); - } - }else + //printf ("staged-copy- read dep signals\n"); + device->_stagingBuffer[1]->CopyDeviceToHost(1,dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); + } + else { - device->_staging_buffer[1]->CopyDeviceToHost(dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); + device->_stagingBuffer[1]->CopyDeviceToHost(0,dst, src, sizeBytes, depSignalCnt ? &depSignal : NULL); } // The copy completes before returning so can reset queue to empty: this->wait(crit, true); @@ -1631,7 +1629,7 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const hsa_agent_t srcAgent = *(static_cast(srcPtrInfo._acc.get_hsa_agent())); ihipSignal_t *ihipSignal = allocSignal(crit); - hsa_signal_t copyCompleteSignal = ihipSignal->_hsa_signal; + hsa_signal_t copyCompleteSignal = ihipSignal->_hsaSignal; hsa_signal_store_relaxed(copyCompleteSignal, 1); void *devPtrDst = dstPtrInfo._devicePointer; @@ -1662,15 +1660,15 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const hsa_agent_t dstAgent = * (static_cast (dstPtrInfo._acc.get_hsa_agent())); hsa_agent_t srcAgent = * (static_cast (srcPtrInfo._acc.get_hsa_agent())); - device->_staging_buffer[1]->CopyPeerToPeer(dst, dstAgent, src, srcAgent, sizeBytes, depSignalCnt ? &depSignal : NULL); + device->_stagingBuffer[1]->CopyPeerToPeer(dst, dstAgent, src, srcAgent, sizeBytes, depSignalCnt ? &depSignal : NULL); // The copy completes before returning so can reset queue to empty: this->wait(crit, true); } else { assert(0); // currently no fallback for this path. - } - + } + } else { // If not special case - these can all be handled by the hsa async copy: ihipCommand_t commandType; @@ -1681,7 +1679,7 @@ void ihipStream_t::copySync(LockedAccessor_StreamCrit_t &crit, void* dst, const // Get a completion signal: ihipSignal_t *ihipSignal = allocSignal(crit); - hsa_signal_t copyCompleteSignal = ihipSignal->_hsa_signal; + hsa_signal_t copyCompleteSignal = ihipSignal->_hsaSignal; hsa_signal_store_relaxed(copyCompleteSignal, 1); @@ -1713,9 +1711,9 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig { LockedAccessor_StreamCrit_t crit(_criticalData); - ihipDevice_t *device = this->getDevice(); + const ihipCtx_t *ctx = this->getCtx(); - if (device == NULL) { + if ((ctx == nullptr) || (ctx->getDevice() == nullptr)) { throw ihipException(hipErrorInvalidDevice); } @@ -1754,7 +1752,7 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig ihipSignal_t *ihip_signal = allocSignal(crit); - hsa_signal_store_relaxed(ihip_signal->_hsa_signal, 1); + hsa_signal_store_relaxed(ihip_signal->_hsaSignal, 1); if(trueAsync == true){ @@ -1766,9 +1764,9 @@ void ihipStream_t::copyAsync(void* dst, const void* src, size_t sizeBytes, unsig hsa_signal_t depSignal; int depSignalCnt = preCopyCommand(crit, ihip_signal, &depSignal, commandType); - tprintf (DB_SYNC, " copy-async, waitFor=%lu completion=#%lu(%lu)\n", depSignalCnt? depSignal.handle:0x0, ihip_signal->_sig_id, ihip_signal->_hsa_signal.handle); + tprintf (DB_SYNC, " copy-async, waitFor=%lu completion=#%lu(%lu)\n", depSignalCnt? depSignal.handle:0x0, ihip_signal->_sigId, ihip_signal->_hsaSignal.handle); - hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, ihip_signal->_hsa_signal); + hsa_status_t hsa_status = hsa_amd_memory_async_copy(dst, dstAgent, src, srcAgent, sizeBytes, depSignalCnt, depSignalCnt ? &depSignal:0x0, ihip_signal->_hsaSignal); if (hsa_status == HSA_STATUS_SUCCESS) { @@ -1799,12 +1797,12 @@ hipError_t hipHccGetAccelerator(int deviceId, hc::accelerator *acc) { HIP_INIT_API(deviceId, acc); - ihipDevice_t *d = ihipGetDevice(deviceId); + const ihipDevice_t *device = ihipGetDevice(deviceId); hipError_t err; - if (d == NULL) { + if (device == NULL) { err = hipErrorInvalidDevice; } else { - *acc = d->_acc; + *acc = device->_acc; err = hipSuccess; } return ihipLogStatus(err); @@ -1820,8 +1818,8 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **a HIP_INIT_API(stream, av); if (stream == hipStreamNull ) { - ihipDevice_t *device = ihipGetTlsDefaultDevice(); - stream = device->_default_stream; + ihipCtx_t *device = ihipGetTlsDefaultCtx(); + stream = device->_defaultStream; } *av = &(stream->_av); @@ -1834,12 +1832,12 @@ hipError_t hipHccGetAcceleratorView(hipStream_t stream, hc::accelerator_view **a // TODO - describe naming convention. ihip _. No accessors. No early returns from functions. Set status to success at top, only set error codes in implementation. No tabs. // Caps convention _ or camelCase // if { } -// Should use ihip* data structures inside code rather than app-facing hip. For example, use ihipDevice_t (rather than hipDevice_t), ihipStream_t (rather than hipStream_t). +// Should use ihip* data structures inside code rather than app-facing hip. For example, use ihipCtx_t (rather than hipDevice_t), ihipStream_t (rather than hipStream_t). // locked_ // TODO - describe MT strategy // //// TODO - add identifier numbers for streams and devices to help with debugging. #if ONE_OBJECT_FILE -#include "staging_buffer.cpp" +#include "unpinned_copy_engine.cpp" #endif diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 94442f4698..c0282dc372 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -1,21 +1,21 @@ /* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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. -*/ + Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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. + */ #include "hip_runtime.h" #include "hcc_detail/hip_hcc.h" @@ -120,18 +120,19 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) hipError_t hip_status = hipSuccess; - auto device = ihipGetTlsDefaultDevice(); + auto ctx = ihipGetTlsDefaultCtx(); - if (device) { + if (ctx) { + auto device = ctx->getWriteableDevice(); const unsigned am_flags = 0; *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags); if (sizeBytes && (*ptr == NULL)) { hip_status = hipErrorMemoryAllocation; } else { - hc::am_memtracker_update(*ptr, device->_device_index, 0); + hc::am_memtracker_update(*ptr, device->_deviceId, 0); { - LockedAccessor_DeviceCrit_t crit(device->criticalData()); + LockedAccessor_CtxCrit_t crit(ctx->criticalData()); if (crit->peerCnt()) { hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); } @@ -152,15 +153,17 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) hipError_t hip_status = hipSuccess; - auto device = ihipGetTlsDefaultDevice(); + auto ctx = ihipGetTlsDefaultCtx(); - if(device){ + if(ctx){ + // am_alloc requires writeable __acc, perhaps could be refactored? + auto device = ctx->getWriteableDevice(); if(flags == hipHostMallocDefault){ *ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned); if(sizeBytes < 1 && (*ptr == NULL)){ hip_status = hipErrorMemoryAllocation; - }else{ - hc::am_memtracker_update(*ptr, device->_device_index, amHostPinned); + } else { + hc::am_memtracker_update(*ptr, device->_deviceId, amHostPinned); } tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr); } else if(flags & hipHostMallocMapped){ @@ -168,9 +171,9 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) if(sizeBytes && (*ptr == NULL)){ hip_status = hipErrorMemoryAllocation; }else{ - hc::am_memtracker_update(*ptr, device->_device_index, flags); + hc::am_memtracker_update(*ptr, device->_deviceId, flags); { - LockedAccessor_DeviceCrit_t crit(device->criticalData()); + LockedAccessor_CtxCrit_t crit(ctx->criticalData()); if (crit->peerCnt()) { hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); } @@ -199,63 +202,67 @@ hipError_t hipMallocHost(void** ptr, size_t sizeBytes) } // width in bytes -hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height) { +hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height) +{ + HIP_INIT_API(ptr, pitch, width, height); - HIP_INIT_API(ptr, pitch, width, height); + hipError_t hip_status = hipSuccess; - hipError_t hip_status = hipSuccess; + if(width == 0 || height == 0) + return ihipLogStatus(hipErrorUnknown); - if(width == 0 || height == 0) - return ihipLogStatus(hipErrorUnknown); + // hardcoded 128 bytes + *pitch = ((((int)width-1)/128) + 1)*128; + const size_t sizeBytes = (*pitch)*height; - // hardcoded 128 bytes - *pitch = ((((int)width-1)/128) + 1)*128; - const size_t sizeBytes = (*pitch)*height; + auto ctx = ihipGetTlsDefaultCtx(); - auto device = ihipGetTlsDefaultDevice(); + //err = hipMalloc(ptr, (*pitch)*height); + if (ctx) { + auto device = ctx->getWriteableDevice(); - //err = hipMalloc(ptr, (*pitch)*height); - if (device) { - const unsigned am_flags = 0; - *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags); + const unsigned am_flags = 0; + *ptr = hc::am_alloc(sizeBytes, device->_acc, am_flags); - if (sizeBytes && (*ptr == NULL)) { - hip_status = hipErrorMemoryAllocation; - } else { - hc::am_memtracker_update(*ptr, device->_device_index, 0); - { - LockedAccessor_DeviceCrit_t crit(device->criticalData()); - if (crit->peerCnt() > 1) { // peerCnt includes self so only call allow_access if other peers involved: - hsa_status_t hsa_status = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); - if (hsa_status != HSA_STATUS_SUCCESS) { + if (sizeBytes && (*ptr == NULL)) { hip_status = hipErrorMemoryAllocation; - } + } else { + hc::am_memtracker_update(*ptr, device->_deviceId, 0); + { + LockedAccessor_CtxCrit_t crit(ctx->criticalData()); + if (crit->peerCnt() > 1) { // peerCnt includes self so only call allow_access if other peers involved: + hsa_status_t hsa_status = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); + if (hsa_status != HSA_STATUS_SUCCESS) { + hip_status = hipErrorMemoryAllocation; + } + } + } } - } + } else { + hip_status = hipErrorMemoryAllocation; } - } else { - hip_status = hipErrorMemoryAllocation; - } - - return ihipLogStatus(hip_status); + return ihipLogStatus(hip_status); } -hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) { - hipChannelFormatDesc cd; - cd.x = x; cd.y = y; cd.z = z; cd.w = w; - cd.f = f; - return cd; + +hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f) +{ + hipChannelFormatDesc cd; + cd.x = x; cd.y = y; cd.z = z; cd.w = w; + cd.f = f; + return cd; } + hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, - size_t width, size_t height, unsigned int flags) { - + size_t width, size_t height, unsigned int flags) +{ HIP_INIT_API(array, desc, width, height, flags); hipError_t hip_status = hipSuccess; - auto device = ihipGetTlsDefaultDevice(); + auto ctx = ihipGetTlsDefaultCtx(); *array = (hipArray*)malloc(sizeof(hipArray)); array[0]->width = width; @@ -265,41 +272,42 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, void ** ptr = &array[0]->data; - if (device) { - const unsigned am_flags = 0; - const size_t size = width*height; + if (ctx) { + auto device = ctx->getWriteableDevice(); + const unsigned am_flags = 0; + const size_t size = width*height; - switch(desc->f) { - case hipChannelFormatKindSigned: - *ptr = hc::am_alloc(size*sizeof(int), device->_acc, am_flags); - break; - case hipChannelFormatKindUnsigned: - *ptr = hc::am_alloc(size*sizeof(unsigned int), device->_acc, am_flags); - break; - case hipChannelFormatKindFloat: - *ptr = hc::am_alloc(size*sizeof(float), device->_acc, am_flags); - break; - case hipChannelFormatKindNone: - *ptr = hc::am_alloc(size*sizeof(size_t), device->_acc, am_flags); - break; - default: - hip_status = hipErrorUnknown; - break; - } - if (size && (*ptr == NULL)) { - hip_status = hipErrorMemoryAllocation; - } else { - hc::am_memtracker_update(*ptr, device->_device_index, 0); - { - LockedAccessor_DeviceCrit_t crit(device->criticalData()); - if (crit->peerCnt() > 1) { // peerCnt includes self so only call allow_access if other peers involved: - hsa_status_t hsa_status = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); - if (hsa_status != HSA_STATUS_SUCCESS) { - hip_status = hipErrorMemoryAllocation; - } - } - } - } + switch(desc->f) { + case hipChannelFormatKindSigned: + *ptr = hc::am_alloc(size*sizeof(int), device->_acc, am_flags); + break; + case hipChannelFormatKindUnsigned: + *ptr = hc::am_alloc(size*sizeof(unsigned int), device->_acc, am_flags); + break; + case hipChannelFormatKindFloat: + *ptr = hc::am_alloc(size*sizeof(float), device->_acc, am_flags); + break; + case hipChannelFormatKindNone: + *ptr = hc::am_alloc(size*sizeof(size_t), device->_acc, am_flags); + break; + default: + hip_status = hipErrorUnknown; + break; + } + if (size && (*ptr == NULL)) { + hip_status = hipErrorMemoryAllocation; + } else { + hc::am_memtracker_update(*ptr, device->_deviceId, 0); + { + LockedAccessor_CtxCrit_t crit(ctx->criticalData()); + if (crit->peerCnt() > 1) { // peerCnt includes self so only call allow_access if other peers involved: + hsa_status_t hsa_status = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr); + if (hsa_status != HSA_STATUS_SUCCESS) { + hip_status = hipErrorMemoryAllocation; + } + } + } + } } else { hip_status = hipErrorMemoryAllocation; @@ -314,24 +322,24 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) { HIP_INIT_API(flagsPtr, hostPtr); - hipError_t hip_status = hipSuccess; + hipError_t hip_status = hipSuccess; - hc::accelerator acc; - hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); - am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr); - if(status == AM_SUCCESS){ - *flagsPtr = amPointerInfo._appAllocationFlags; - if(*flagsPtr == 0){ - hip_status = hipErrorInvalidValue; - } - else{ - hip_status = hipSuccess; - } - tprintf(DB_MEM, " %s: host ptr=%p\n", __func__, hostPtr); - }else{ - hip_status = hipErrorInvalidValue; - } - return ihipLogStatus(hip_status); + hc::accelerator acc; + hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); + am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, hostPtr); + if(status == AM_SUCCESS){ + *flagsPtr = amPointerInfo._appAllocationFlags; + if(*flagsPtr == 0){ + hip_status = hipErrorInvalidValue; + } + else{ + hip_status = hipSuccess; + } + tprintf(DB_MEM, " %s: host ptr=%p\n", __func__, hostPtr); + }else{ + hip_status = hipErrorInvalidValue; + } + return ihipLogStatus(hip_status); } @@ -342,7 +350,7 @@ hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags) hipError_t hip_status = hipSuccess; - auto device = ihipGetTlsDefaultDevice(); + auto ctx = ihipGetTlsDefaultCtx(); if(hostPtr == NULL){ return ihipLogStatus(hipErrorInvalidValue); } @@ -353,24 +361,25 @@ hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags) if(am_status == AM_SUCCESS){ hip_status = hipErrorHostMemoryAlreadyRegistered; - }else{ - auto device = ihipGetTlsDefaultDevice(); + } else { + auto ctx = ihipGetTlsDefaultCtx(); if(hostPtr == NULL){ return ihipLogStatus(hipErrorInvalidValue); } - if(device){ + if (ctx) { + auto device = ctx->getWriteableDevice(); if(flags == hipHostRegisterDefault || flags == hipHostRegisterPortable || flags == hipHostRegisterMapped){ std::vectorvecAcc; for(int i=0;i_acc); } am_status = hc::am_memory_host_lock(device->_acc, hostPtr, sizeBytes, &vecAcc[0], vecAcc.size()); if(am_status == AM_SUCCESS){ hip_status = hipSuccess; - }else{ + } else { hip_status = hipErrorMemoryAllocation; } - }else{ + } else { hip_status = hipErrorInvalidValue; } } @@ -382,11 +391,12 @@ hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags) hipError_t hipHostUnregister(void *hostPtr) { HIP_INIT_API(hostPtr); - auto device = ihipGetTlsDefaultDevice(); + auto ctx = ihipGetTlsDefaultCtx(); hipError_t hip_status = hipSuccess; if(hostPtr == NULL){ hip_status = hipErrorInvalidValue; }else{ + auto device = ctx->getWriteableDevice(); am_status_t am_status = hc::am_memory_host_unlock(device->_acc, hostPtr); if(am_status != AM_SUCCESS){ hip_status = hipErrorHostMemoryNotRegistered; @@ -402,17 +412,17 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t cou HIP_INIT_API(symbolName, src, count, offset, kind); #ifdef USE_MEMCPYTOSYMBOL - if(kind != hipMemcpyHostToDevice) - { - return ihipLogStatus(hipErrorInvalidValue); - } - auto device = ihipGetTlsDefaultDevice(); + if(kind != hipMemcpyHostToDevice) + { + return ihipLogStatus(hipErrorInvalidValue); + } + auto ctx = ihipGetTlsDefaultCtx(); //hsa_signal_t depSignal; - //int depSignalCnt = device._default_stream->preCopyCommand(NULL, &depSignal, ihipCommandCopyH2D); + //int depSignalCnt = ctx._default_stream->preCopyCommand(NULL, &depSignal, ihipCommandCopyH2D); assert(0); // Need to properly synchronize the copy - do something with depSignal if != NULL. - device->_acc.memcpy_symbol(symbolName, (void*) src,count, offset); + ctx->_acc.memcpy_symbol(symbolName, (void*) src,count, offset); #endif return ihipLogStatus(hipSuccess); } @@ -476,110 +486,183 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp // dpitch, spitch, and width in bytes hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, - size_t width, size_t height, hipMemcpyKind kind) { + size_t width, size_t height, hipMemcpyKind kind) { - HIP_INIT_API(dst, dpitch, src, spitch, width, height, kind); + HIP_INIT_API(dst, dpitch, src, spitch, width, height, kind); - if(width > dpitch || width > spitch) - return ihipLogStatus(hipErrorUnknown); + if(width > dpitch || width > spitch) + return ihipLogStatus(hipErrorUnknown); - hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); - hc::completion_future marker; + hc::completion_future marker; - hipError_t e = hipSuccess; + hipError_t e = hipSuccess; - try { - for(int i = 0; i < height; ++i) { - stream->locked_copySync((unsigned char*)dst + i*dpitch, (unsigned char*)src + i*spitch, width, kind); + try { + for(int i = 0; i < height; ++i) { + stream->locked_copySync((unsigned char*)dst + i*dpitch, (unsigned char*)src + i*spitch, width, kind); + } + } + catch (ihipException ex) { + e = ex._code; } - } - catch (ihipException ex) { - e = ex._code; - } - return ihipLogStatus(e); + return ihipLogStatus(e); } // wOffset, width, and spitch in bytes hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, - size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { + size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { - HIP_INIT_API(dst, wOffset, hOffset, src, spitch, width, height, kind); + HIP_INIT_API(dst, wOffset, hOffset, src, spitch, width, height, kind); - hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); - hc::completion_future marker; + hc::completion_future marker; - hipError_t e = hipSuccess; + hipError_t e = hipSuccess; - size_t byteSize; - if(dst) { - switch(dst[0].f) { - case hipChannelFormatKindSigned: - byteSize = sizeof(int); - break; - case hipChannelFormatKindUnsigned: - byteSize = sizeof(unsigned int); - break; - case hipChannelFormatKindFloat: - byteSize = sizeof(float); - break; - case hipChannelFormatKindNone: - byteSize = sizeof(size_t); - break; - default: - byteSize = 0; - break; + size_t byteSize; + if(dst) { + switch(dst[0].f) { + case hipChannelFormatKindSigned: + byteSize = sizeof(int); + break; + case hipChannelFormatKindUnsigned: + byteSize = sizeof(unsigned int); + break; + case hipChannelFormatKindFloat: + byteSize = sizeof(float); + break; + case hipChannelFormatKindNone: + byteSize = sizeof(size_t); + break; + default: + byteSize = 0; + break; + } + } else { + return ihipLogStatus(hipErrorUnknown); } - } else { - return ihipLogStatus(hipErrorUnknown); - } - if((wOffset + width > (dst->width * byteSize)) || width > spitch) { - return ihipLogStatus(hipErrorUnknown); - } - - size_t src_w = spitch; - size_t dst_w = (dst->width)*byteSize; - - try { - for(int i = 0; i < height; ++i) { - stream->locked_copySync((unsigned char*)dst->data + i*dst_w, (unsigned char*)src + i*src_w, width, kind); + if((wOffset + width > (dst->width * byteSize)) || width > spitch) { + return ihipLogStatus(hipErrorUnknown); } - } - catch (ihipException ex) { - e = ex._code; - } - return ihipLogStatus(e); + size_t src_w = spitch; + size_t dst_w = (dst->width)*byteSize; + + try { + for(int i = 0; i < height; ++i) { + stream->locked_copySync((unsigned char*)dst->data + i*dst_w, (unsigned char*)src + i*src_w, width, kind); + } + } + catch (ihipException ex) { + e = ex._code; + } + + return ihipLogStatus(e); } hipError_t hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, - const void* src, size_t count, hipMemcpyKind kind) { + const void* src, size_t count, hipMemcpyKind kind) { - HIP_INIT_API(dst, wOffset, hOffset, src, count, kind); + HIP_INIT_API(dst, wOffset, hOffset, src, count, kind); - hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); - hc::completion_future marker; + hc::completion_future marker; - hipError_t e = hipSuccess; + hipError_t e = hipSuccess; - try { - stream->locked_copySync((char *)dst->data + wOffset, src, count, kind); - } - catch (ihipException ex) { - e = ex._code; - } + try { + stream->locked_copySync((char *)dst->data + wOffset, src, count, kind); + } + catch (ihipException ex) { + e = ex._code; + } - return ihipLogStatus(e); + return ihipLogStatus(e); } + +template +hc::completion_future +ihipMemsetKernel(hipStream_t stream, T * ptr, T val, size_t sizeBytes) +{ + int wg = std::min((unsigned)8, stream->getDevice()->_computeUnits); + const int threads_per_wg = 256; + + int threads = wg * threads_per_wg; + if (threads > sizeBytes) { + threads = ((sizeBytes + threads_per_wg - 1) / threads_per_wg) * threads_per_wg; + } + + + hc::extent<1> ext(threads); + auto ext_tile = ext.tile(threads_per_wg); + + hc::completion_future cf = + hc::parallel_for_each( + stream->_av, + ext_tile, + [=] (hc::tiled_index<1> idx) + __attribute__((hc)) + { + int offset = amp_get_global_id(0); + // TODO-HCC - change to hc_get_local_size() + int stride = amp_get_local_size(0) * hc_get_num_groups(0) ; + + for (int i=offset; i +hc::completion_future +ihipMemcpyKernel(hipStream_t stream, T * c, const T * a, size_t sizeBytes) +{ + int wg = std::min((unsigned)8, stream->getDevice()->_computeUnits); + const int threads_per_wg = 256; + + int threads = wg * threads_per_wg; + if (threads > sizeBytes) { + threads = ((sizeBytes + threads_per_wg - 1) / threads_per_wg) * threads_per_wg; + } + + + hc::extent<1> ext(threads); + auto ext_tile = ext.tile(threads_per_wg); + + hc::completion_future cf = + hc::parallel_for_each( + stream->_av, + ext_tile, + [=] (hc::tiled_index<1> idx) + __attribute__((hc)) + { + int offset = amp_get_global_id(0); + // TODO-HCC - change to hc_get_local_size() + int stride = amp_get_local_size(0) * hc_get_num_groups(0) ; + + for (int i=offset; igetWriteableDevice(); if (total) { - *total = hipDevice->_props.totalGlobalMem; + *total = device->_props.totalGlobalMem; } if (free) { // TODO - replace with kernel-level for reporting free memory: size_t deviceMemSize, hostMemSize, userMemSize; - hc::am_memtracker_sizeinfo(hipDevice->_acc, &deviceMemSize, &hostMemSize, &userMemSize); + hc::am_memtracker_sizeinfo(device->_acc, &deviceMemSize, &hostMemSize, &userMemSize); printf ("deviceMemSize=%zu\n", deviceMemSize); - *free = hipDevice->_props.totalGlobalMem - deviceMemSize; + *free = device->_props.totalGlobalMem - deviceMemSize; } } else { @@ -722,8 +806,8 @@ hipError_t hipFree(void* ptr) hipError_t hipStatus = hipErrorInvalidDevicePointer; - // Synchronize to ensure all work has finished. - ihipGetTlsDefaultDevice()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish. + // Synchronize to ensure all work has finished. + ihipGetTlsDefaultCtx()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish. if (ptr) { hc::accelerator acc; @@ -748,8 +832,8 @@ hipError_t hipHostFree(void* ptr) { HIP_INIT_API(ptr); - // Synchronize to ensure all work has finished. - ihipGetTlsDefaultDevice()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish. + // Synchronize to ensure all work has finished. + ihipGetTlsDefaultCtx()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish. hipError_t hipStatus = hipErrorInvalidValue; @@ -780,26 +864,26 @@ hipError_t hipFreeHost(void* ptr) hipError_t hipFreeArray(hipArray* array) { - HIP_INIT_API(array); + HIP_INIT_API(array); - hipError_t hipStatus = hipErrorInvalidDevicePointer; + hipError_t hipStatus = hipErrorInvalidDevicePointer; - // Synchronize to ensure all work has finished. - ihipGetTlsDefaultDevice()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish. + // Synchronize to ensure all work has finished. + ihipGetTlsDefaultCtx()->locked_waitAllStreams(); // ignores non-blocking streams, this waits for all activity to finish. - if(array->data) { - hc::accelerator acc; - hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); - am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, array->data); - if(status == AM_SUCCESS){ - if(amPointerInfo._hostPointer == NULL){ - hc::am_free(array->data); - hipStatus = hipSuccess; - } + if(array->data) { + hc::accelerator acc; + hc::AmPointerInfo amPointerInfo(NULL, NULL, 0, acc, 0, 0); + am_status_t status = hc::am_memtracker_getinfo(&amPointerInfo, array->data); + if(status == AM_SUCCESS){ + if(amPointerInfo._hostPointer == NULL){ + hc::am_free(array->data); + hipStatus = hipSuccess; + } + } } - } - return ihipLogStatus(hipStatus); + return ihipLogStatus(hipStatus); } // Stubs of threadfence operations diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp new file mode 100644 index 0000000000..1421eb0329 --- /dev/null +++ b/projects/hip/src/hip_module.cpp @@ -0,0 +1,254 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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. +*/ + +#include "hip_runtime.h" +#include "hsa/hsa.h" +#include "hsa/hsa_ext_amd.h" +#include "hcc_detail/hip_hcc.h" +#include "hcc_detail/trace_helper.h" +#include + +//TODO Use Pool APIs from HCC to get memory regions. + +namespace hipdrv{ +hsa_status_t findSystemRegions(hsa_region_t region, void *data){ + hsa_region_segment_t segment_id; + hsa_region_get_info(region, HSA_REGION_INFO_SEGMENT, &segment_id); + + if(segment_id != HSA_REGION_SEGMENT_GLOBAL){ + return HSA_STATUS_SUCCESS; + } + + hsa_region_global_flag_t flags; + hsa_region_get_info(region, HSA_REGION_INFO_GLOBAL_FLAGS, &flags); + + hsa_region_t *reg = (hsa_region_t*)data; + + if(flags & HSA_REGION_GLOBAL_FLAG_FINE_GRAINED){ + *reg = region; + } + + return HSA_STATUS_SUCCESS; +} + +hsa_status_t findKernArgRegions(hsa_region_t region, void *data){ + hsa_region_segment_t segment_id; + hsa_region_get_info(region, HSA_REGION_INFO_SEGMENT, &segment_id); + + if(segment_id != HSA_REGION_SEGMENT_GLOBAL){ + return HSA_STATUS_SUCCESS; + } + + hsa_region_global_flag_t flags; + hsa_region_get_info(region, HSA_REGION_INFO_GLOBAL_FLAGS, &flags); + + hsa_region_t *reg = (hsa_region_t*)data; + + if(flags & HSA_REGION_GLOBAL_FLAG_KERNARG){ + *reg = region; + } + + return HSA_STATUS_SUCCESS; +} + + +} + + + +hipError_t hipModuleLoad(hipModule *module, const char *fname){ + HIP_INIT_API(fname); + hipError_t ret = hipSuccess; + if(module == NULL){ + ret = hipErrorInvalidValue; + } + auto ctx = ihipGetTlsDefaultCtx(); + if(ctx == nullptr){ + ret = hipErrorInvalidContext; + }else{ + int deviceId = ctx->getDevice()->_deviceId; + ihipDevice_t *currentDevice = ihipGetDevice(deviceId); + std::ifstream in(fname, std::ios::binary | std::ios::ate); + if(!in){ + return hipErrorFileNotFound; + }else{ + size_t size = std::string::size_type(in.tellg()); + void *p = NULL; + hsa_agent_t agent = currentDevice->_hsaAgent; + hsa_region_t sysRegion; + hsa_status_t status = hsa_agent_iterate_regions(agent, hipdrv::findSystemRegions, &sysRegion); + status = hsa_memory_allocate(sysRegion, size, (void**)&p); + if(status != HSA_STATUS_SUCCESS){ + return hipErrorOutOfMemory; + } + char *ptr = (char*)p; + if(!ptr){ + return hipErrorOutOfMemory; + std::cout<<"Error: failed to allocate memory for code object"<(in), + std::istreambuf_iterator(), ptr); + status = hsa_code_object_deserialize(ptr, size, NULL, &obj); + if(status != HSA_STATUS_SUCCESS){ + return hipErrorSharedObjectInitFailed; + } + *module = obj.handle; + } + } + return ret; +} + +hipError_t hipModuleGetFunction(hipFunction *func, hipModule hmod, const char *name){ + HIP_INIT_API(name); + auto ctx = ihipGetTlsDefaultCtx(); + hipError_t ret = hipSuccess; + if(name == nullptr || hmod == 0){ + return hipErrorInvalidValue; + } + if(ctx == nullptr){ + ret = hipErrorInvalidContext; + }else{ + int deviceId = ctx->getDevice()->_deviceId; + ihipDevice_t *currentDevice = ihipGetDevice(deviceId); + hsa_agent_t gpuAgent = (hsa_agent_t)currentDevice->_hsaAgent; + + hsa_status_t status; + hsa_executable_symbol_t kernel_symbol; + hsa_executable_t executable; + status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &executable); + if(status != HSA_STATUS_SUCCESS){ + return hipErrorNotInitialized; + } + hsa_code_object_t obj; + obj.handle = hmod; + status = hsa_executable_load_code_object(executable, gpuAgent, obj, NULL); + if(status != HSA_STATUS_SUCCESS){ + return hipErrorNotInitialized; + } + status = hsa_executable_freeze(executable, NULL); + status = hsa_executable_get_symbol(executable, NULL, name, gpuAgent, 0, &kernel_symbol); + if(status != HSA_STATUS_SUCCESS){ + return hipErrorNotFound; + } + status = hsa_executable_symbol_get_info(kernel_symbol, + HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, + func); + if(status != HSA_STATUS_SUCCESS){ + return hipErrorNotFound; + } + } + return ret; +} + +hipError_t hipLaunchModuleKernel(hipFunction f, + uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ, + uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ, + uint32_t sharedMemBytes, hipStream_t hStream, + void **kernelParams, void **extra){ + HIP_INIT_API(f); + auto ctx = ihipGetTlsDefaultCtx(); + hipError_t ret = hipSuccess; + if(ctx == nullptr){ + ret = hipErrorInvalidDevice; + }else{ + int deviceId = ctx->getDevice()->_deviceId; + ihipDevice_t *currentDevice = ihipGetDevice(deviceId); + hsa_agent_t gpuAgent = (hsa_agent_t)currentDevice->_hsaAgent; + + void *config[5] = {0}; + size_t kernSize; + + if(extra != NULL){ + memcpy(config, extra, sizeof(size_t)*5); + if(config[0] == HIP_LAUNCH_PARAM_BUFFER_POINTER && config[2] == HIP_LAUNCH_PARAM_BUFFER_SIZE && config[4] == HIP_LAUNCH_PARAM_END){ + kernSize = *(size_t*)(config[3]); + }else{ + return hipErrorNotInitialized; + } + }else{ + return hipErrorInvalidValue; + } +/* +Kernel argument preparation. +*/ + + hsa_region_t kernArg; + hsa_status_t status = hsa_agent_iterate_regions(gpuAgent, hipdrv::findKernArgRegions, &kernArg); + void *kern; + status = hsa_memory_allocate(kernArg, kernSize, &kern); + if(status != HSA_STATUS_SUCCESS){ + return hipErrorLaunchOutOfResources; + } + memcpy(kern, config[1], kernSize); + + +/* +Pre kernel launch + + stream = ihipSyncAndResolveStream(stream); + stream->lockopen_preKernelCommand(); + hc::accelerator_view av = &stream->_av; + hc::completion_future cf = new hc::completion_future; +*/ + + hStream = ihipSyncAndResolveStream(hStream); + hc::accelerator_view *av = &hStream->_av; + hsa_queue_t *Queue = (hsa_queue_t*)av->get_hsa_queue(); + hsa_signal_t signal; + status = hsa_signal_create(1, 0, NULL, &signal); + +/* +Creating the packets +*/ + + const uint32_t queue_mask = Queue->size-1; + uint32_t packet_index = hsa_queue_load_write_index_relaxed(Queue); + hsa_kernel_dispatch_packet_t *dispatch_packet = &(((hsa_kernel_dispatch_packet_t*)(Queue->base_address))[packet_index & queue_mask]); + + dispatch_packet->completion_signal = signal; + dispatch_packet->workgroup_size_x = blockDimX; + dispatch_packet->workgroup_size_y = blockDimY; + dispatch_packet->workgroup_size_z = blockDimZ; + dispatch_packet->grid_size_x = blockDimX * gridDimX; + dispatch_packet->grid_size_y = blockDimY * gridDimY; + dispatch_packet->grid_size_z = blockDimZ * gridDimZ; + + dispatch_packet->group_segment_size = 0; + dispatch_packet->private_segment_size = sharedMemBytes; + dispatch_packet->kernarg_address = kern; + dispatch_packet->kernel_object = f; + uint16_t header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | + (1 << HSA_PACKET_HEADER_BARRIER) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE); + + + uint16_t setup = 1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS; + uint32_t header32 = header | (setup << 16); + + __atomic_store_n((uint32_t*)(dispatch_packet), header32, __ATOMIC_RELEASE); + + hsa_queue_store_write_index_relaxed(Queue, packet_index+1); + hsa_signal_store_relaxed(Queue->doorbell_signal, packet_index); + hsa_signal_value_t value = hsa_signal_wait_acquire(signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); + } + return ret; +} diff --git a/projects/hip/src/hip_peer.cpp b/projects/hip/src/hip_peer.cpp index cec8017b0c..fd51599815 100644 --- a/projects/hip/src/hip_peer.cpp +++ b/projects/hip/src/hip_peer.cpp @@ -23,25 +23,32 @@ THE SOFTWARE. #include "hcc_detail/hip_hcc.h" #include "hcc_detail/trace_helper.h" + +// Peer access functions. +// There are two flavors: +// - one where contexts are specified with hipCtx_t type. +// - one where contexts are specified with integer deviceIds, that are mapped to the primary context for that device. +// The implementation contains a set of internal ihip* functions which operate on contexts. Then the +// public APIs are thin wrappers which call into this internal implementations. +// TODO - actually not yet - currently the integer deviceId flavors just call the context APIs. need to fix. + /** * HCC returns 0 in *canAccessPeer ; Need to update this function when RT supports P2P */ //--- -hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDeviceId) +hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t thisCtx, hipCtx_t peerCtx) { - HIP_INIT_API(canAccessPeer, deviceId, peerDeviceId); + HIP_INIT_API(canAccessPeer, thisCtx, peerCtx); hipError_t err = hipSuccess; - auto thisDevice = ihipGetDevice(deviceId); - auto peerDevice = ihipGetDevice(peerDeviceId); - if ((thisDevice != NULL) && (peerDevice != NULL)) { - if (deviceId == peerDeviceId) { + if ((thisCtx != NULL) && (peerCtx != NULL)) { + if (thisCtx == peerCtx) { *canAccessPeer = 0; } else { #if USE_PEER_TO_PEER>=2 - *canAccessPeer = peerDevice->_acc.get_is_peer(thisDevice->_acc); + *canAccessPeer = peerCtx->getDevice()->_acc.get_is_peer(thisCtx->getDevice()->_acc); #else *canAccessPeer = 0; #endif @@ -60,33 +67,32 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDe //--- // Disable visibility of this device into memory allocated on peer device. // Remove this device from peer device peerlist. -hipError_t hipDeviceDisablePeerAccess (int peerDeviceId) +hipError_t hipDeviceDisablePeerAccess (hipCtx_t peerCtx) { - HIP_INIT_API(peerDeviceId); + HIP_INIT_API(peerCtx); hipError_t err = hipSuccess; - auto thisDevice = ihipGetTlsDefaultDevice(); - auto peerDevice = ihipGetDevice(peerDeviceId); - if ((thisDevice != NULL) && (peerDevice != NULL)) { + auto thisCtx = ihipGetTlsDefaultCtx(); + if ((thisCtx != NULL) && (peerCtx != NULL)) { #if USE_PEER_TO_PEER>=2 - // Return true if thisDevice can access peerDevice's memory: - bool canAccessPeer = peerDevice->_acc.get_is_peer(thisDevice->_acc); + // Return true if thisCtx can access peerCtx's memory: + bool canAccessPeer = peerCtx->getDevice()->_acc.get_is_peer(thisCtx->getDevice()->_acc); #else bool canAccessPeer = 0; #endif if (! canAccessPeer) { err = hipErrorInvalidDevice; // P2P not allowed between these devices. - } else if (thisDevice == peerDevice) { + } else if (thisCtx == peerCtx) { err = hipErrorInvalidDevice; // Can't disable peer access to self. } else { - LockedAccessor_DeviceCrit_t peerCrit(peerDevice->criticalData()); - bool changed = peerCrit->removePeer(thisDevice); + LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData()); + bool changed = peerCrit->removePeer(thisCtx); if (changed) { #if USE_PEER_TO_PEER>=3 // Update the peers for all memory already saved in the tracker: - am_memtracker_update_peers(peerDevice->_acc, peerCrit->peerCnt(), peerCrit->peerAgents()); + am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(), peerCrit->peerAgents()); #endif } else { err = hipErrorPeerAccessNotEnabled; // never enabled P2P access. @@ -103,24 +109,23 @@ hipError_t hipDeviceDisablePeerAccess (int peerDeviceId) //--- // Allow the current device to see all memory allocated on peerDevice. // This should add this device to the peer-device peer list. -hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags) +hipError_t hipDeviceEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags) { - HIP_INIT_API(peerDeviceId, flags); + HIP_INIT_API(peerCtx, flags); hipError_t err = hipSuccess; if (flags != 0) { err = hipErrorInvalidValue; } else { - auto thisDevice = ihipGetTlsDefaultDevice(); - auto peerDevice = ihipGetDevice(peerDeviceId); - if (thisDevice == peerDevice) { + auto thisCtx = ihipGetTlsDefaultCtx(); + if (thisCtx == peerCtx) { err = hipErrorInvalidDevice; // Can't enable peer access to self. - } else if ((thisDevice != NULL) && (peerDevice != NULL)) { - LockedAccessor_DeviceCrit_t peerCrit(peerDevice->criticalData()); - bool isNewPeer = peerCrit->addPeer(thisDevice); + } else if ((thisCtx != NULL) && (peerCtx != NULL)) { + LockedAccessor_CtxCrit_t peerCrit(peerCtx->criticalData()); + bool isNewPeer = peerCrit->addPeer(thisCtx); if (isNewPeer) { #if USE_PEER_TO_PEER>=3 - am_memtracker_update_peers(peerDevice->_acc, peerCrit->peerCnt(), peerCrit->peerAgents()); + am_memtracker_update_peers(peerCtx->getDevice()->_acc, peerCrit->peerCnt(), peerCrit->peerAgents()); #endif } else { err = hipErrorPeerAccessAlreadyEnabled; @@ -135,20 +140,17 @@ hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags) //--- -hipError_t hipMemcpyPeer (void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes) +hipError_t hipMemcpyPeer (void* dst, hipCtx_t dstCtx, const void* src, hipCtx_t srcCtx, size_t sizeBytes) { - HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes); + HIP_INIT_API(dst, dstCtx, src, srcCtx, sizeBytes); // HCC has a unified memory architecture so device specifiers are not required. return hipMemcpy(dst, src, sizeBytes, hipMemcpyDefault); }; -/** - * This function uses a synchronous copy - */ //--- -hipError_t hipMemcpyPeerAsync (void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream) +hipError_t hipMemcpyPeerAsync (void* dst, hipCtx_t dstDevice, const void* src, hipCtx_t srcDevice, size_t sizeBytes, hipStream_t stream) { HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes, stream); // HCC has a unified memory architecture so device specifiers are not required. @@ -156,19 +158,48 @@ hipError_t hipMemcpyPeerAsync (void* dst, int dstDevice, const void* src, int }; -/** - * @return #hipSuccess - */ -//--- -hipError_t hipDriverGetVersion(int *driverVersion) + +//============================================================================= +// These are the flavors that accept integer deviceIDs. +// Implementations map these to primary contexts and call the internal functions above. +//============================================================================= + +hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDeviceId) { - HIP_INIT_API(driverVersion); - - if (driverVersion) { - *driverVersion = 4; - } - - return ihipLogStatus(hipSuccess); + HIP_INIT_API(canAccessPeer, deviceId, peerDeviceId); + return hipDeviceCanAccessPeer(canAccessPeer, ihipGetPrimaryCtx(deviceId), ihipGetPrimaryCtx(peerDeviceId)); } +hipError_t hipDeviceDisablePeerAccess (int peerDeviceId) +{ + HIP_INIT_API(peerDeviceId); + + return hipDeviceDisablePeerAccess(ihipGetPrimaryCtx(peerDeviceId)); +} + + +hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags) +{ + HIP_INIT_API(peerDeviceId, flags); + + return hipDeviceEnablePeerAccess(ihipGetPrimaryCtx(peerDeviceId), flags); +} + + +hipError_t hipMemcpyPeer (void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes) +{ + HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes); + return hipMemcpyPeer(dst, ihipGetPrimaryCtx(dstDevice), src, ihipGetPrimaryCtx(srcDevice), sizeBytes); +} + + +hipError_t hipMemcpyPeerAsync (void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream) +{ + HIP_INIT_API(dst, dstDevice, src, srcDevice, sizeBytes, stream); + return hipMemcpyPeerAsync(dst, ihipGetPrimaryCtx(dstDevice), src, ihipGetPrimaryCtx(srcDevice), sizeBytes, stream); +} + + + + diff --git a/projects/hip/src/hip_stream.cpp b/projects/hip/src/hip_stream.cpp index d62abc49e2..a204e2f79a 100644 --- a/projects/hip/src/hip_stream.cpp +++ b/projects/hip/src/hip_stream.cpp @@ -30,23 +30,29 @@ THE SOFTWARE. //--- hipError_t ihipStreamCreate(hipStream_t *stream, unsigned int flags) { - ihipDevice_t *device = ihipGetTlsDefaultDevice(); - hc::accelerator acc = device->_acc; + ihipCtx_t *ctx = ihipGetTlsDefaultCtx(); - // TODO - se try-catch loop to detect memory exception? - // - // - //Note this is an execute_in_order queue, so all kernels submitted will atuomatically wait for prev to complete: - //This matches CUDA stream behavior: + hipError_t e = hipSuccess; - auto istream = new ihipStream_t(device->_device_index, acc.create_view(), flags); + if (ctx) { + hc::accelerator acc = ctx->getWriteableDevice()->_acc; - device->locked_addStream(istream); + // TODO - se try-catch loop to detect memory exception? + // + //Note this is an execute_in_order queue, so all kernels submitted will atuomatically wait for prev to complete: + //This matches CUDA stream behavior: - *stream = istream; - tprintf(DB_SYNC, "hipStreamCreate, stream=%p\n", *stream); + auto istream = new ihipStream_t(ctx, acc.create_view(), flags); - return hipSuccess; + ctx->locked_addStream(istream); + + *stream = istream; + tprintf(DB_SYNC, "hipStreamCreate, stream=%p\n", *stream); + } else { + e = hipErrorInvalidDevice; + } + + return ihipLogStatus(e); } @@ -98,8 +104,8 @@ hipError_t hipStreamSynchronize(hipStream_t stream) hipError_t e = hipSuccess; if (stream == NULL) { - ihipDevice_t *device = ihipGetTlsDefaultDevice(); - device->locked_syncDefaultStream(true/*waitOnSelf*/); + ihipCtx_t *ctx = ihipGetTlsDefaultCtx(); + ctx->locked_syncDefaultStream(true/*waitOnSelf*/); } else { stream->locked_wait(); e = hipSuccess; @@ -122,17 +128,17 @@ hipError_t hipStreamDestroy(hipStream_t stream) //--- Drain the stream: if (stream == NULL) { - ihipDevice_t *device = ihipGetTlsDefaultDevice(); - device->locked_syncDefaultStream(true/*waitOnSelf*/); + ihipCtx_t *ctx = ihipGetTlsDefaultCtx(); + ctx->locked_syncDefaultStream(true/*waitOnSelf*/); } else { stream->locked_wait(); e = hipSuccess; } - ihipDevice_t *device = stream->getDevice(); + ihipCtx_t *ctx = stream->getCtx(); - if (device) { - device->locked_removeStream(stream); + if (ctx) { + ctx->locked_removeStream(stream); delete stream; } else { e = hipErrorInvalidResourceHandle; diff --git a/projects/hip/src/staging_buffer.cpp b/projects/hip/src/unpinned_copy_engine.cpp similarity index 55% rename from projects/hip/src/staging_buffer.cpp rename to projects/hip/src/unpinned_copy_engine.cpp index c6c23089bd..5501c66f9d 100644 --- a/projects/hip/src/staging_buffer.cpp +++ b/projects/hip/src/unpinned_copy_engine.cpp @@ -21,44 +21,83 @@ THE SOFTWARE. #include "hsa_ext_amd.h" -#include "hcc_detail/staging_buffer.h" +#include "hcc_detail/unpinned_copy_engine.h" #ifdef HIP_HCC #include "hcc_detail/hip_runtime.h" #include "hcc_detail/hip_hcc.h" #define THROW_ERROR(e) throw ihipException(e) #else -#define THROW_ERROR(e) throw -#define tprintf(trace_level, ...) +#define THROW_ERROR(e) throw +#define tprintf(trace_level, ...) #endif -extern hsa_agent_t g_cpu_agent; // defined in hip_hcc.cpp +void errorCheck(hsa_status_t hsa_error_code, int line_num, std::string str) { + if ((hsa_error_code != HSA_STATUS_SUCCESS)&& (hsa_error_code != HSA_STATUS_INFO_BREAK)) { + printf("HSA reported error!\n In file: %s\nAt line: %d\n", str.c_str(),line_num); + } +} + +#define ErrorCheck(x) errorCheck(x, __LINE__, __FILE__) +hsa_amd_memory_pool_t sys_pool_; + +hsa_status_t findGlobalPool(hsa_amd_memory_pool_t pool, void* data) { + if (NULL == data) { + return HSA_STATUS_ERROR_INVALID_ARGUMENT; + } + + hsa_status_t err; + hsa_amd_segment_t segment; + uint32_t flag; + err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_SEGMENT, &segment); + ErrorCheck(err); + + err = hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &flag); + ErrorCheck(err); + if ((HSA_AMD_SEGMENT_GLOBAL == segment) && + (flag & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED)) { + *((hsa_amd_memory_pool_t*)data) = pool; + } + return HSA_STATUS_SUCCESS; +} //------------------------------------------------------------------------------------------------- -StagingBuffer::StagingBuffer(hsa_agent_t hsaAgent, hsa_region_t systemRegion, size_t bufferSize, int numBuffers) : - _hsa_agent(hsaAgent), +UnpinnedCopyEngine::UnpinnedCopyEngine(hsa_agent_t hsaAgent, hsa_agent_t cpuAgent, size_t bufferSize, int numBuffers, int thresholdH2DDirectStaging,int thresholdH2DStagingPinInPlace,int thresholdD2H) : + _hsaAgent(hsaAgent), + _cpuAgent(cpuAgent), _bufferSize(bufferSize), - _numBuffers(numBuffers > _max_buffers ? _max_buffers : numBuffers) + _numBuffers(numBuffers > _max_buffers ? _max_buffers : numBuffers), + _hipH2DTransferThresholdDirectOrStaging(thresholdH2DDirectStaging), + _hipH2DTransferThresholdStagingOrPininplace(thresholdH2DStagingPinInPlace), + _hipD2HTransferThreshold(thresholdD2H) { + hsa_status_t err = hsa_amd_agent_iterate_memory_pools(_cpuAgent, findGlobalPool, &sys_pool_); + ErrorCheck(err); for (int i=0; i<_numBuffers; i++) { // TODO - experiment with alignment here. - hsa_status_t s1 = hsa_memory_allocate(systemRegion, _bufferSize, (void**) (&_pinnedStagingBuffer[i]) ); + err = hsa_amd_memory_pool_allocate(sys_pool_, _bufferSize, 0, (void**)(&_pinnedStagingBuffer[i])); + ErrorCheck(err); - if ((s1 != HSA_STATUS_SUCCESS) || (_pinnedStagingBuffer[i] == NULL)) { + if ((err != HSA_STATUS_SUCCESS) || (_pinnedStagingBuffer[i] == NULL)) { THROW_ERROR(hipErrorMemoryAllocation); } + + err = hsa_amd_agents_allow_access(1, &hsaAgent, NULL, _pinnedStagingBuffer[i]); + ErrorCheck(err); + hsa_signal_create(0, 0, NULL, &_completion_signal[i]); hsa_signal_create(0, 0, NULL, &_completion_signal2[i]); } + }; //--- -StagingBuffer::~StagingBuffer() +UnpinnedCopyEngine::~UnpinnedCopyEngine() { for (int i=0; i<_numBuffers; i++) { if (_pinnedStagingBuffer[i]) { - hsa_memory_free(_pinnedStagingBuffer[i]); + hsa_amd_memory_pool_free(_pinnedStagingBuffer[i]); _pinnedStagingBuffer[i] = NULL; } hsa_signal_destroy(_completion_signal[i]); @@ -71,9 +110,9 @@ StagingBuffer::~StagingBuffer() //--- //Copies sizeBytes from src to dst, using either a copy to a staging buffer or a staged pin-in-place strategy //IN: dst - dest pointer - must be accessible from host CPU. -//IN: src - src pointer for copy. Must be accessible from agent this buffer is associated with (via _hsa_agent) +//IN: src - src pointer for copy. Must be accessible from agent this buffer is associated with (via _hsaAgent) //IN: waitFor - hsaSignal to wait for - the copy will begin only when the specified dependency is resolved. May be NULL indicating no dependency. -void StagingBuffer::CopyHostToDevicePinInPlace(void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor) +void UnpinnedCopyEngine::CopyHostToDevicePinInPlace(void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor) { std::lock_guard l (_copy_lock); @@ -88,19 +127,15 @@ void StagingBuffer::CopyHostToDevicePinInPlace(void* dst, const void* src, size_ THROW_ERROR (hipErrorInvalidValue); } int bufferIndex = 0; -#if 0 - for (int64_t bytesRemaining=sizeBytes; bytesRemaining>0 ; bytesRemaining -= _bufferSize) { - size_t theseBytes = (bytesRemaining > _bufferSize) ? _bufferSize : bytesRemaining; -#endif size_t theseBytes= sizeBytes; //tprintf (DB_COPY2, "H2D: waiting... on completion signal handle=%lu\n", _completion_signal[bufferIndex].handle); //hsa_signal_wait_acquire(_completion_signal[bufferIndex], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); //void * masked_srcp = (void*) ((uintptr_t)srcp & (uintptr_t)(~0x3f)) ; // TODO void *locked_srcp; - //hsa_status_t hsa_status = hsa_amd_memory_lock(masked_srcp, theseBytes, &_hsa_agent, 1, &locked_srcp); - hsa_status_t hsa_status = hsa_amd_memory_lock(const_cast (srcp), theseBytes, &_hsa_agent, 1, &locked_srcp); + //hsa_status_t hsa_status = hsa_amd_memory_lock(masked_srcp, theseBytes, &_hsaAgent, 1, &locked_srcp); + hsa_status_t hsa_status = hsa_amd_memory_lock(const_cast (srcp), theseBytes, &_hsaAgent, 1, &locked_srcp); //tprintf (DB_COPY2, "H2D: bytesRemaining=%zu: pin-in-place:%p+%zu bufferIndex[%d]\n", bytesRemaining, srcp, theseBytes, bufferIndex); //printf ("status=%x srcp=%p, masked_srcp=%p, locked_srcp=%p\n", hsa_status, srcp, masked_srcp, locked_srcp); @@ -110,7 +145,7 @@ void StagingBuffer::CopyHostToDevicePinInPlace(void* dst, const void* src, size_ hsa_signal_store_relaxed(_completion_signal[bufferIndex], 1); - hsa_status = hsa_amd_memory_async_copy(dstp, _hsa_agent, locked_srcp, g_cpu_agent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]); + hsa_status = hsa_amd_memory_async_copy(dstp, _hsaAgent, locked_srcp, _cpuAgent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]); //tprintf (DB_COPY2, "H2D: bytesRemaining=%zu: async_copy %zu bytes %p to %p status=%x\n", bytesRemaining, theseBytes, _pinnedStagingBuffer[bufferIndex], dstp, hsa_status); if (hsa_status != HSA_STATUS_SUCCESS) { @@ -119,26 +154,8 @@ void StagingBuffer::CopyHostToDevicePinInPlace(void* dst, const void* src, size_ tprintf (DB_COPY2, "H2D: waiting... on completion signal handle=%lu\n", _completion_signal[bufferIndex].handle); hsa_signal_wait_acquire(_completion_signal[bufferIndex], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); hsa_amd_memory_unlock(const_cast (srcp)); -#if 0 - srcp += theseBytes; - dstp += theseBytes; - if (++bufferIndex >= _numBuffers) { - bufferIndex = 0; - } -#endif - // Assume subsequent commands are dependent on previous and don't need dependency after first copy submitted, HIP_ONESHOT_COPY_DEP=1 - waitFor = NULL; -#if 0 -// } - - // TODO - - printf ("unpin the memory\n"); - - - for (int i=0; i<_numBuffers; i++) { - hsa_signal_wait_acquire(_completion_signal[i], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); - } -#endif + // Assume subsequent commands are dependent on previous and don't need dependency after first copy submitted, HIP_ONESHOT_COPY_DEP=1 + waitFor = NULL; } @@ -147,62 +164,70 @@ void StagingBuffer::CopyHostToDevicePinInPlace(void* dst, const void* src, size_ //--- //Copies sizeBytes from src to dst, using either a copy to a staging buffer or a staged pin-in-place strategy //IN: dst - dest pointer - must be accessible from host CPU. -//IN: src - src pointer for copy. Must be accessible from agent this buffer is associated with (via _hsa_agent) +//IN: src - src pointer for copy. Must be accessible from agent this buffer is associated with (via _hsaAgent) //IN: waitFor - hsaSignal to wait for - the copy will begin only when the specified dependency is resolved. May be NULL indicating no dependency. -void StagingBuffer::CopyHostToDevice(void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor) +void UnpinnedCopyEngine::CopyHostToDevice(int tempIndex,int isLargeBar,void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor) { - std::lock_guard l (_copy_lock); - - const char *srcp = static_cast (src); - char *dstp = static_cast (dst); - - for (int i=0; i<_numBuffers; i++) { - hsa_signal_store_relaxed(_completion_signal[i], 0); + if((tempIndex==1)&&(isLargeBar)&&(sizeBytes < _hipH2DTransferThresholdDirectOrStaging)){ + memcpy(dst,src,sizeBytes); + std::atomic_thread_fence(std::memory_order_release); + } + else if((tempIndex==1) && (sizeBytes > _hipH2DTransferThresholdStagingOrPininplace)){ + CopyHostToDevicePinInPlace(dst, src, sizeBytes, waitFor); } + else + { + std::lock_guard l (_copy_lock); - if (sizeBytes >= UINT64_MAX/2) { - THROW_ERROR (hipErrorInvalidValue); - } - int bufferIndex = 0; - for (int64_t bytesRemaining=sizeBytes; bytesRemaining>0 ; bytesRemaining -= _bufferSize) { + const char *srcp = static_cast (src); + char *dstp = static_cast (dst); - size_t theseBytes = (bytesRemaining > _bufferSize) ? _bufferSize : bytesRemaining; - - tprintf (DB_COPY2, "H2D: waiting... on completion signal handle=%lu\n", _completion_signal[bufferIndex].handle); - hsa_signal_wait_acquire(_completion_signal[bufferIndex], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); - - tprintf (DB_COPY2, "H2D: bytesRemaining=%zu: copy %zu bytes %p to stagingBuf[%d]:%p\n", bytesRemaining, theseBytes, srcp, bufferIndex, _pinnedStagingBuffer[bufferIndex]); - // TODO - use uncached memcpy, someday. - memcpy(_pinnedStagingBuffer[bufferIndex], srcp, theseBytes); - - - hsa_signal_store_relaxed(_completion_signal[bufferIndex], 1); - - hsa_status_t hsa_status = hsa_amd_memory_async_copy(dstp, _hsa_agent, _pinnedStagingBuffer[bufferIndex], g_cpu_agent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]); - tprintf (DB_COPY2, "H2D: bytesRemaining=%zu: async_copy %zu bytes %p to %p status=%x\n", bytesRemaining, theseBytes, _pinnedStagingBuffer[bufferIndex], dstp, hsa_status); - - if (hsa_status != HSA_STATUS_SUCCESS) { - THROW_ERROR ((hipErrorRuntimeMemory)); + for (int i=0; i<_numBuffers; i++) { + hsa_signal_store_relaxed(_completion_signal[i], 0); } - srcp += theseBytes; - dstp += theseBytes; - if (++bufferIndex >= _numBuffers) { - bufferIndex = 0; + if (sizeBytes >= UINT64_MAX/2) { + THROW_ERROR (hipErrorInvalidValue); + } + int bufferIndex = 0; + for (int64_t bytesRemaining=sizeBytes; bytesRemaining>0 ; bytesRemaining -= _bufferSize) { + + size_t theseBytes = (bytesRemaining > _bufferSize) ? _bufferSize : bytesRemaining; + + tprintf (DB_COPY2, "H2D: waiting... on completion signal handle=%lu\n", _completion_signal[bufferIndex].handle); + hsa_signal_wait_acquire(_completion_signal[bufferIndex], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); + + tprintf (DB_COPY2, "H2D: bytesRemaining=%zu: copy %zu bytes %p to stagingBuf[%d]:%p\n", bytesRemaining, theseBytes, srcp, bufferIndex, _pinnedStagingBuffer[bufferIndex]); + // TODO - use uncached memcpy, someday. + memcpy(_pinnedStagingBuffer[bufferIndex], srcp, theseBytes); + + + hsa_signal_store_relaxed(_completion_signal[bufferIndex], 1); + hsa_status_t hsa_status = hsa_amd_memory_async_copy(dstp, _hsaAgent, _pinnedStagingBuffer[bufferIndex], _cpuAgent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]); + tprintf (DB_COPY2, "H2D: bytesRemaining=%zu: async_copy %zu bytes %p to %p status=%x\n", bytesRemaining, theseBytes, _pinnedStagingBuffer[bufferIndex], dstp, hsa_status); + if (hsa_status != HSA_STATUS_SUCCESS) { + THROW_ERROR ((hipErrorRuntimeMemory)); + } + + srcp += theseBytes; + dstp += theseBytes; + if (++bufferIndex >= _numBuffers) { + bufferIndex = 0; + } + + // Assume subsequent commands are dependent on previous and don't need dependency after first copy submitted, HIP_ONESHOT_COPY_DEP=1 + waitFor = NULL; } - // Assume subsequent commands are dependent on previous and don't need dependency after first copy submitted, HIP_ONESHOT_COPY_DEP=1 - waitFor = NULL; - } - - for (int i=0; i<_numBuffers; i++) { - hsa_signal_wait_acquire(_completion_signal[i], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); - } + for (int i=0; i<_numBuffers; i++) { + hsa_signal_wait_acquire(_completion_signal[i], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); + } + } } -void StagingBuffer::CopyDeviceToHostPinInPlace(void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor) +void UnpinnedCopyEngine::CopyDeviceToHostPinInPlace(void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor) { std::lock_guard l (_copy_lock); @@ -220,7 +245,7 @@ void StagingBuffer::CopyDeviceToHostPinInPlace(void* dst, const void* src, size_ size_t theseBytes= sizeBytes; void *locked_destp; - hsa_status_t hsa_status = hsa_amd_memory_lock(const_cast (dstp), theseBytes, &_hsa_agent, 1, &locked_destp); + hsa_status_t hsa_status = hsa_amd_memory_lock(const_cast (dstp), theseBytes, &_hsaAgent, 1, &locked_destp); if (hsa_status != HSA_STATUS_SUCCESS) { @@ -229,7 +254,7 @@ void StagingBuffer::CopyDeviceToHostPinInPlace(void* dst, const void* src, size_ hsa_signal_store_relaxed(_completion_signal[bufferIndex], 1); - hsa_status = hsa_amd_memory_async_copy(locked_destp,g_cpu_agent , srcp, _hsa_agent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]); + hsa_status = hsa_amd_memory_async_copy(locked_destp,_cpuAgent , srcp, _hsaAgent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]); if (hsa_status != HSA_STATUS_SUCCESS) { THROW_ERROR (hipErrorRuntimeMemory); @@ -244,70 +269,77 @@ void StagingBuffer::CopyDeviceToHostPinInPlace(void* dst, const void* src, size_ //--- //Copies sizeBytes from src to dst, using either a copy to a staging buffer or a staged pin-in-place strategy -//IN: dst - dest pointer - must be accessible from agent this buffer is associated with (via _hsa_agent). +//IN: dst - dest pointer - must be accessible from agent this buffer is associated with (via _hsaAgent). //IN: src - src pointer for copy. Must be accessible from host CPU. //IN: waitFor - hsaSignal to wait for - the copy will begin only when the specified dependency is resolved. May be NULL indicating no dependency. -void StagingBuffer::CopyDeviceToHost(void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor) +void UnpinnedCopyEngine::CopyDeviceToHost(int tempIndex,void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor) { - std::lock_guard l (_copy_lock); - - const char *srcp0 = static_cast (src); - char *dstp1 = static_cast (dst); - - for (int i=0; i<_numBuffers; i++) { - hsa_signal_store_relaxed(_completion_signal[i], 0); + if((tempIndex==1) && (sizeBytes> _hipD2HTransferThreshold)){ + CopyDeviceToHostPinInPlace(dst, src, sizeBytes, waitFor); } + else + { + std::lock_guard l (_copy_lock); - if (sizeBytes >= UINT64_MAX/2) { - THROW_ERROR (hipErrorInvalidValue); - } + const char *srcp0 = static_cast (src); + char *dstp1 = static_cast (dst); - int64_t bytesRemaining0 = sizeBytes; // bytes to copy from dest into staging buffer. - int64_t bytesRemaining1 = sizeBytes; // bytes to copy from staging buffer into final dest + for (int i=0; i<_numBuffers; i++) { + hsa_signal_store_relaxed(_completion_signal[i], 0); + } - while (bytesRemaining1 > 0) { - // First launch the async copies to copy from dest to host - for (int bufferIndex = 0; (bytesRemaining0>0) && (bufferIndex < _numBuffers); bytesRemaining0 -= _bufferSize, bufferIndex++) { + if (sizeBytes >= UINT64_MAX/2) { + THROW_ERROR (hipErrorInvalidValue); + } - size_t theseBytes = (bytesRemaining0 > _bufferSize) ? _bufferSize : bytesRemaining0; + int64_t bytesRemaining0 = sizeBytes; // bytes to copy from dest into staging buffer. + int64_t bytesRemaining1 = sizeBytes; // bytes to copy from staging buffer into final dest - tprintf (DB_COPY2, "D2H: bytesRemaining0=%zu async_copy %zu bytes src:%p to staging:%p\n", bytesRemaining0, theseBytes, srcp0, _pinnedStagingBuffer[bufferIndex]); - hsa_signal_store_relaxed(_completion_signal[bufferIndex], 1); - hsa_status_t hsa_status = hsa_amd_memory_async_copy(_pinnedStagingBuffer[bufferIndex], g_cpu_agent, srcp0, _hsa_agent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]); - if (hsa_status != HSA_STATUS_SUCCESS) { - THROW_ERROR (hipErrorRuntimeMemory); + while (bytesRemaining1 > 0) + { + // First launch the async copies to copy from dest to host + for (int bufferIndex = 0; (bytesRemaining0>0) && (bufferIndex < _numBuffers); bytesRemaining0 -= _bufferSize, bufferIndex++) { + + size_t theseBytes = (bytesRemaining0 > _bufferSize) ? _bufferSize : bytesRemaining0; + + tprintf (DB_COPY2, "D2H: bytesRemaining0=%zu async_copy %zu bytes src:%p to staging:%p\n", bytesRemaining0, theseBytes, srcp0, _pinnedStagingBuffer[bufferIndex]); + hsa_signal_store_relaxed(_completion_signal[bufferIndex], 1); + hsa_status_t hsa_status = hsa_amd_memory_async_copy(_pinnedStagingBuffer[bufferIndex], _cpuAgent, srcp0, _hsaAgent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]); + if (hsa_status != HSA_STATUS_SUCCESS) { + THROW_ERROR (hipErrorRuntimeMemory); + } + + srcp0 += theseBytes; + + + // Assume subsequent commands are dependent on previous and don't need dependency after first copy submitted, HIP_ONESHOT_COPY_DEP=1 + waitFor = NULL; } - srcp0 += theseBytes; + // Now unload the staging buffers: + for (int bufferIndex=0; (bytesRemaining1>0) && (bufferIndex < _numBuffers); bytesRemaining1 -= _bufferSize, bufferIndex++) { + size_t theseBytes = (bytesRemaining1 > _bufferSize) ? _bufferSize : bytesRemaining1; - // Assume subsequent commands are dependent on previous and don't need dependency after first copy submitted, HIP_ONESHOT_COPY_DEP=1 - waitFor = NULL; - } + tprintf (DB_COPY2, "D2H: wait_completion[%d] bytesRemaining=%zu\n", bufferIndex, bytesRemaining1); + hsa_signal_wait_acquire(_completion_signal[bufferIndex], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); - // Now unload the staging buffers: - for (int bufferIndex=0; (bytesRemaining1>0) && (bufferIndex < _numBuffers); bytesRemaining1 -= _bufferSize, bufferIndex++) { + tprintf (DB_COPY2, "D2H: bytesRemaining1=%zu copy %zu bytes stagingBuf[%d]:%p to dst:%p\n", bytesRemaining1, theseBytes, bufferIndex, _pinnedStagingBuffer[bufferIndex], dstp1); + memcpy(dstp1, _pinnedStagingBuffer[bufferIndex], theseBytes); - size_t theseBytes = (bytesRemaining1 > _bufferSize) ? _bufferSize : bytesRemaining1; - - tprintf (DB_COPY2, "D2H: wait_completion[%d] bytesRemaining=%zu\n", bufferIndex, bytesRemaining1); - hsa_signal_wait_acquire(_completion_signal[bufferIndex], HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_ACTIVE); - - tprintf (DB_COPY2, "D2H: bytesRemaining1=%zu copy %zu bytes stagingBuf[%d]:%p to dst:%p\n", bytesRemaining1, theseBytes, bufferIndex, _pinnedStagingBuffer[bufferIndex], dstp1); - memcpy(dstp1, _pinnedStagingBuffer[bufferIndex], theseBytes); - - dstp1 += theseBytes; - } + dstp1 += theseBytes; + } + } } } //--- //Copies sizeBytes from src to dst, using either a copy to a staging buffer or a staged pin-in-place strategy -//IN: dst - dest pointer - must be accessible from agent this buffer is associated with (via _hsa_agent). +//IN: dst - dest pointer - must be accessible from agent this buffer is associated with (via _hsaAgent). //IN: src - src pointer for copy. Must be accessible from host CPU. //IN: waitFor - hsaSignal to wait for - the copy will begin only when the specified dependency is resolved. May be NULL indicating no dependency. -void StagingBuffer::CopyPeerToPeer(void* dst, hsa_agent_t dstAgent, const void* src, hsa_agent_t srcAgent, size_t sizeBytes, hsa_signal_t *waitFor) +void UnpinnedCopyEngine::CopyPeerToPeer(void* dst, hsa_agent_t dstAgent, const void* src, hsa_agent_t srcAgent, size_t sizeBytes, hsa_signal_t *waitFor) { std::lock_guard l (_copy_lock); @@ -337,7 +369,7 @@ void StagingBuffer::CopyPeerToPeer(void* dst, hsa_agent_t dstAgent, const void* tprintf (DB_COPY2, "P2P: bytesRemaining0=%zu async_copy %zu bytes src:%p to staging:%p\n", bytesRemaining0, theseBytes, srcp0, _pinnedStagingBuffer[bufferIndex]); hsa_signal_store_relaxed(_completion_signal[bufferIndex], 1); - hsa_status_t hsa_status = hsa_amd_memory_async_copy(_pinnedStagingBuffer[bufferIndex], g_cpu_agent, srcp0, srcAgent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]); + hsa_status_t hsa_status = hsa_amd_memory_async_copy(_pinnedStagingBuffer[bufferIndex], _cpuAgent, srcp0, srcAgent, theseBytes, waitFor ? 1:0, waitFor, _completion_signal[bufferIndex]); if (hsa_status != HSA_STATUS_SUCCESS) { THROW_ERROR (hipErrorRuntimeMemory); } @@ -345,8 +377,8 @@ void StagingBuffer::CopyPeerToPeer(void* dst, hsa_agent_t dstAgent, const void* srcp0 += theseBytes; - // Assume subsequent commands are dependent on previous and don't need dependency after first copy submitted, HIP_ONESHOT_COPY_DEP=1 - waitFor = NULL; + // Assume subsequent commands are dependent on previous and don't need dependency after first copy submitted, HIP_ONESHOT_COPY_DEP=1 + waitFor = NULL; } // Now unload the staging buffers: @@ -365,8 +397,8 @@ void StagingBuffer::CopyPeerToPeer(void* dst, hsa_agent_t dstAgent, const void* tprintf (DB_COPY2, "P2P: bytesRemaining1=%zu copy %zu bytes stagingBuf[%d]:%p to device:%p\n", bytesRemaining1, theseBytes, bufferIndex, _pinnedStagingBuffer[bufferIndex], dstp1); hsa_signal_store_relaxed(_completion_signal2[bufferIndex], 1); - hsa_status_t hsa_status = hsa_amd_memory_async_copy(dstp1, dstAgent, _pinnedStagingBuffer[bufferIndex], g_cpu_agent /*not used*/, theseBytes, - hostWait ? 0:1, hostWait ? NULL : &_completion_signal[bufferIndex], + hsa_status_t hsa_status = hsa_amd_memory_async_copy(dstp1, dstAgent, _pinnedStagingBuffer[bufferIndex], _cpuAgent /*not used*/, theseBytes, + hostWait ? 0:1, hostWait ? NULL : &_completion_signal[bufferIndex], _completion_signal2[bufferIndex]); dstp1 += theseBytes; diff --git a/projects/hip/tests/src/CMakeLists.txt b/projects/hip/tests/src/CMakeLists.txt index 127ebb2507..86ba712ead 100644 --- a/projects/hip/tests/src/CMakeLists.txt +++ b/projects/hip/tests/src/CMakeLists.txt @@ -1,5 +1,7 @@ cmake_minimum_required (VERSION 2.6) +# remove CMAKE_CXX_COMPILER entry from cache since it will be pointing to hipcc +unset(CMAKE_CXX_COMPILER CACHE) message (CMAKE_CXX_COMPILER = ${CMAKE_CXX_COMPILER} ) project (HIP_Unit_Tests) @@ -175,7 +177,7 @@ build_hip_executable (hipFuncSetDevice hipFuncSetDevice.cpp) build_hip_executable (hipFuncDeviceSynchronize hipFuncDeviceSynchronize.cpp) build_hip_executable (hipPeerToPeer_simple hipPeerToPeer_simple.cpp) build_hip_executable (hipTestMemcpyPin hipTestMemcpyPin.cpp) -#build_hip_executable (hipDynamicShared hipDynamicShared.cpp) +build_hip_executable (hipDynamicShared hipDynamicShared.cpp) build_hip_executable (hipLaunchParm hipLaunchParm.cpp) if (${HIP_PLATFORM} STREQUAL "hcc") @@ -214,9 +216,10 @@ endif() make_hipify_test(specialFunc.cu ) -#make_test(hipDynamicShared " ") +make_test(hipDynamicShared " ") # Add subdirs here: +add_subdirectory(context) add_subdirectory(deviceLib) add_subdirectory(runtimeApi) add_subdirectory(kernel) diff --git a/projects/hip/tests/src/context/CMakeLists.txt b/projects/hip/tests/src/context/CMakeLists.txt new file mode 100644 index 0000000000..d04985d001 --- /dev/null +++ b/projects/hip/tests/src/context/CMakeLists.txt @@ -0,0 +1,9 @@ +cmake_minimum_required (VERSION 2.6) + +# Functions for kernel attributes (grid_launch, __launch_bounds__, etc) +project (kernel) + +include_directories( ${HIPTEST_SOURCE_DIR} ) + +build_hip_executable_libcpp (hipCtx_simple hipCtx_simple.cpp) +make_test(hipCtx_simple " " ) diff --git a/projects/hip/tests/src/context/hipCtx_simple.cpp b/projects/hip/tests/src/context/hipCtx_simple.cpp new file mode 100644 index 0000000000..882cf44f6d --- /dev/null +++ b/projects/hip/tests/src/context/hipCtx_simple.cpp @@ -0,0 +1,47 @@ +/* +Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "hip_runtime.h" +#include "test_common.h" + +int main(int argc, char *argv[]) +{ + HipTest::parseStandardArguments(argc, argv, true); + + HIPCHECK(hipInit(0)); + + hipDevice_t device; + hipDevice_t device1; + hipCtx_t ctx; + hipCtx_t ctx1; + + HIPCHECK(hipDeviceGetFromId(&device, 0)); + HIPCHECK(hipCtxCreate(&ctx, 0, device)); + HIPCHECK(hipCtxGetCurrent(&ctx1)); + HIPCHECK(hipCtxGetDevice(&device1)); + HIPCHECK(hipCtxPopCurrent(&ctx1)); + HIPCHECK(hipCtxGetCurrent(&ctx1)); + + HIPCHECK(hipCtxDestroy(ctx)); + + passed(); +}; diff --git a/projects/hip/tests/src/deviceLib/hipDoublePrecisionMathDevice.cpp b/projects/hip/tests/src/deviceLib/hipDoublePrecisionMathDevice.cpp index 90c8ae1aa6..40343053f5 100644 --- a/projects/hip/tests/src/deviceLib/hipDoublePrecisionMathDevice.cpp +++ b/projects/hip/tests/src/deviceLib/hipDoublePrecisionMathDevice.cpp @@ -47,9 +47,9 @@ __device__ void double_precision_math_functions() //cyl_bessel_i1(0.0); erf(0.0); erfc(0.0); - //erfcinv(2.0); - //erfcx(0.0); - //erfinv(1.0); + erfcinv(2.0); + erfcx(0.0); + erfinv(1.0); exp(0.0); exp10(0.0); exp2(0.0); @@ -67,46 +67,46 @@ __device__ void double_precision_math_functions() isfinite(0.0); isinf(0.0); isnan(0.0); - //j0(0.0); - //j1(0.0); - //jn(-1.0, 1.0); + j0(0.0); + j1(0.0); + jn(-1.0, 1.0); ldexp(0.0, 0); //lgamma(1.0); - //llrint(0.0); - //llround(0.0); + llrint(0.0); + llround(0.0); log(1.0); log10(1.0); log1p(-1.0); log2(1.0); logb(1.0); - //lrint(0.0); - //lround(0.0); + lrint(0.0); + lround(0.0); //modf(0.0, &fX); nan("1"); nearbyint(0.0); //nextafter(0.0); //fX = 1.0; norm(1, &fX); - //norm3d(1.0, 0.0, 0.0); - //norm4d(1.0, 0.0, 0.0, 0.0); - //normcdf(0.0); + norm3d(1.0, 0.0, 0.0); + norm4d(1.0, 0.0, 0.0, 0.0); + normcdf(0.0); //normcdfinv(1.0); pow(1.0, 0.0); - //rcbrt(1.0); + rcbrt(1.0); remainder(2.0, 1.0); //remquo(1.0, 2.0, &iX); - //rhypot(0.0, 1.0); - //rint(1.0); - //fX = 1.0; rnorm(1, &fX); - //rnorm3d(0.0, 0.0, 1.0); - //rnorm4d(0.0, 0.0, 0.0, 1.0); + rhypot(0.0, 1.0); + rint(1.0); + fX = 1.0; rnorm(1, &fX); + rnorm3d(0.0, 0.0, 1.0); + rnorm4d(0.0, 0.0, 0.0, 1.0); round(0.0); rsqrt(1.0); - //scalbln(0.0, 1); + scalbln(0.0, 1); scalbn(0.0, 1); signbit(1.0); sin(0.0); - //sincos(0.0, &fX, &fY); - //sincospi(0.0, &fX, &fY); + sincos(0.0, &fX, &fY); + sincospi(0.0, &fX, &fY); sinh(0.0); sinpi(0.0); sqrt(0.0); @@ -114,9 +114,9 @@ __device__ void double_precision_math_functions() tanh(0.0); tgamma(2.0); trunc(0.0); - //y0(1.0); - //y1(1.0); - //yn(1, 1.0); + y0(1.0); + y1(1.0); + yn(1, 1.0); } __global__ void compileDoublePrecisionMathOnDevice(hipLaunchParm lp, int ignored) diff --git a/projects/hip/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp b/projects/hip/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp index 461dc5a609..94fe912e08 100644 --- a/projects/hip/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp +++ b/projects/hip/tests/src/deviceLib/hipDoublePrecisionMathHost.cpp @@ -47,9 +47,9 @@ __host__ void double_precision_math_functions() //cyl_bessel_i1(0.0); erf(0.0); erfc(0.0); - //erfcinv(2.0); - //erfcx(0.0); - //erfinv(1.0); + erfcinv(2.0); + erfcx(0.0); + erfinv(1.0); exp(0.0); exp10(0.0); exp2(0.0); @@ -67,46 +67,46 @@ __host__ void double_precision_math_functions() isfinite(0.0); isinf(0.0); isnan(0.0); - ///j0(0.0); - ///j1(0.0); - ///jn(-1.0, 1.0); + j0(0.0); + j1(0.0); + jn(-1.0, 1.0); ldexp(0.0, 0); - ///lgamma(1.0); - ///llrint(0.0); - ///llround(0.0); + lgamma(1.0); + llrint(0.0); + llround(0.0); log(1.0); log10(1.0); log1p(-1.0); log2(1.0); logb(1.0); - ///lrint(0.0); - ///lround(0.0); + lrint(0.0); + lround(0.0); modf(0.0, &fX); - ///nan("1"); + nan("1"); nearbyint(0.0); //nextafter(0.0); //fX = 1.0; norm(1, &fX); - //norm3d(1.0, 0.0, 0.0); - //norm4d(1.0, 0.0, 0.0, 0.0); - //normcdf(0.0); - //normcdfinv(1.0); + norm3d(1.0, 0.0, 0.0); + norm4d(1.0, 0.0, 0.0, 0.0); + normcdf(0.0); + normcdfinv(1.0); pow(1.0, 0.0); - //rcbrt(1.0); + rcbrt(1.0); remainder(2.0, 1.0); remquo(1.0, 2.0, &iX); - //rhypot(0.0, 1.0); - ///rint(1.0); - //fX = 1.0; rnorm(1, &fX); - //rnorm3d(0.0, 0.0, 1.0); - //rnorm4d(0.0, 0.0, 0.0, 1.0); + rhypot(0.0, 1.0); + rint(1.0); + fX = 1.0; rnorm(1, &fX); + rnorm3d(0.0, 0.0, 1.0); + rnorm4d(0.0, 0.0, 0.0, 1.0); round(0.0); rsqrt(1.0); - ///scalbln(0.0, 1); + scalbln(0.0, 1); scalbn(0.0, 1); signbit(1.0); sin(0.0); sincos(0.0, &fX, &fY); - //sincospi(0.0, &fX, &fY); + sincospi(0.0, &fX, &fY); sinh(0.0); sinpi(0.0); sqrt(0.0); @@ -114,9 +114,9 @@ __host__ void double_precision_math_functions() tanh(0.0); tgamma(2.0); trunc(0.0); - ///y0(1.0); - ///y1(1.0); - ///yn(1, 1.0); + y0(1.0); + y1(1.0); + yn(1, 1.0); } static void compileOnHost() diff --git a/projects/hip/tests/src/deviceLib/hipSinglePrecisionMathDevice.cpp b/projects/hip/tests/src/deviceLib/hipSinglePrecisionMathDevice.cpp index d97cd03dc9..6aa0171ac9 100644 --- a/projects/hip/tests/src/deviceLib/hipSinglePrecisionMathDevice.cpp +++ b/projects/hip/tests/src/deviceLib/hipSinglePrecisionMathDevice.cpp @@ -46,17 +46,17 @@ __device__ void single_precision_math_functions() //cyl_bessel_i0f(0.0f); //cyl_bessel_i1f(0.0f); erfcf(0.0f); - //erfcinvf(2.0f); - //erfcxf(0.0f); + erfcinvf(2.0f); + erfcxf(0.0f); erff(0.0f); - //erfinvf(1.0f); + erfinvf(1.0f); exp10f(0.0f); exp2f(0.0f); expf(0.0f); expm1f(0.0f); fabsf(1.0f); fdimf(1.0f, 0.0f); - //fdividef(0.0f, 1.0f); + fdividef(0.0f, 1.0f); floorf(0.0f); fmaf(1.0f, 2.0f, 3.0f); fmaxf(0.0f, 0.0f); @@ -68,45 +68,45 @@ __device__ void single_precision_math_functions() isfinite(0.0f); isinf(0.0f); isnan(0.0f); - //j0f(0.0f); - //j1f(0.0f); - //jnf(-1.0f, 1.0f); + j0f(0.0f); + j1f(0.0f); + jnf(-1.0f, 1.0f); ldexpf(0.0f, 0); //lgammaf(1.0f); - //llrintf(0.0f); - //llroundf(0.0f); + llrintf(0.0f); + llroundf(0.0f); log10f(1.0f); log1pf(-1.0f); log2f(1.0f); logbf(1.0f); logf(1.0f); - //lrintf(0.0f); - //lroundf(0.0f); + lrintf(0.0f); + lroundf(0.0f); //modff(0.0f, &fX); nanf("1"); nearbyintf(0.0f); //nextafterf(0.0f); - //norm3df(1.0f, 0.0f, 0.0f); - //norm4df(1.0f, 0.0f, 0.0f, 0.0f); - //normcdff(0.0f); - //normcdfinvf(1.0f); - //fX = 1.0f; normf(1, &fX); + norm3df(1.0f, 0.0f, 0.0f); + norm4df(1.0f, 0.0f, 0.0f, 0.0f); + normcdff(0.0f); + normcdfinvf(1.0f); + fX = 1.0f; normf(1, &fX); powf(1.0f, 0.0f); - //rcbrtf(1.0f); + rcbrtf(1.0f); remainderf(2.0f, 1.0f); //remquof(1.0f, 2.0f, &iX); - //rhypotf(0.0f, 1.0f); - //rintf(1.0f); - //rnorm3df(0.0f, 0.0f, 1.0f); - //rnorm4df(0.0f, 0.0f, 0.0f, 1.0f); - //fX = 1.0f; rnormf(1, &fX); + rhypotf(0.0f, 1.0f); + rintf(1.0f); + rnorm3df(0.0f, 0.0f, 1.0f); + rnorm4df(0.0f, 0.0f, 0.0f, 1.0f); + fX = 1.0f; rnormf(1, &fX); roundf(0.0f); rsqrtf(1.0f); - //scalblnf(0.0f, 1); + scalblnf(0.0f, 1); scalbnf(0.0f, 1); signbit(1.0f); - //sincosf(0.0f, &fX, &fY); - //sincospif(0.0f, &fX, &fY); + sincosf(0.0f, &fX, &fY); + sincospif(0.0f, &fX, &fY); sinf(0.0f); sinhf(0.0f); sinpif(0.0f); @@ -115,9 +115,9 @@ __device__ void single_precision_math_functions() tanhf(0.0f); tgammaf(2.0f); truncf(0.0f); - //y0f(1.0f); - //y1f(1.0f); - //ynf(1, 1.0f); + y0f(1.0f); + y1f(1.0f); + ynf(1, 1.0f); } __global__ void compileSinglePrecisionMathOnDevice(hipLaunchParm lp, int ignored) diff --git a/projects/hip/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp b/projects/hip/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp index 67261a0735..8a95bcaad2 100644 --- a/projects/hip/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp +++ b/projects/hip/tests/src/deviceLib/hipSinglePrecisionMathHost.cpp @@ -46,17 +46,17 @@ __host__ void single_precision_math_functions() //cyl_bessel_i0f(0.0f); //cyl_bessel_i1f(0.0f); erfcf(0.0f); - //erfcinvf(2.0f); - //erfcxf(0.0f); + erfcinvf(2.0f); + erfcxf(0.0f); erff(0.0f); - //erfinvf(1.0f); + erfinvf(1.0f); exp10f(0.0f); exp2f(0.0f); expf(0.0f); expm1f(0.0f); fabsf(1.0f); fdimf(1.0f, 0.0f); - //fdividef(0.0f, 1.0f); + fdividef(0.0f, 1.0f); floorf(0.0f); fmaf(1.0f, 2.0f, 3.0f); fmaxf(0.0f, 0.0f); @@ -68,45 +68,45 @@ __host__ void single_precision_math_functions() isfinite(0.0f); isinf(0.0f); isnan(0.0f); - ///j0f(0.0f); - ///j1f(0.0f); - ///jnf(-1.0f, 1.0f); + j0f(0.0f); + j1f(0.0f); + jnf(-1.0f, 1.0f); ldexpf(0.0f, 0); - ///lgammaf(1.0f); - ///llrintf(0.0f); - ///llroundf(0.0f); + lgammaf(1.0f); + llrintf(0.0f); + llroundf(0.0f); log10f(1.0f); log1pf(-1.0f); log2f(1.0f); logbf(1.0f); logf(1.0f); - ///lrintf(0.0f); - ///lroundf(0.0f); + lrintf(0.0f); + lroundf(0.0f); modff(0.0f, &fX); - ///nanf("1"); + nanf("1"); nearbyintf(0.0f); //nextafterf(0.0f); - //norm3df(1.0f, 0.0f, 0.0f); - //norm4df(1.0f, 0.0f, 0.0f, 0.0f); - //normcdff(0.0f); - //normcdfinvf(1.0f); + norm3df(1.0f, 0.0f, 0.0f); + norm4df(1.0f, 0.0f, 0.0f, 0.0f); + normcdff(0.0f); + normcdfinvf(1.0f); //fX = 1.0f; normf(1, &fX); powf(1.0f, 0.0f); - //rcbrtf(1.0f); + rcbrtf(1.0f); remainderf(2.0f, 1.0f); remquof(1.0f, 2.0f, &iX); - //rhypotf(0.0f, 1.0f); - ///rintf(1.0f); - //rnorm3df(0.0f, 0.0f, 1.0f); - //rnorm4df(0.0f, 0.0f, 0.0f, 1.0f); - //fX = 1.0f; rnormf(1, &fX); + rhypotf(0.0f, 1.0f); + rintf(1.0f); + rnorm3df(0.0f, 0.0f, 1.0f); + rnorm4df(0.0f, 0.0f, 0.0f, 1.0f); + fX = 1.0f; rnormf(1, &fX); roundf(0.0f); rsqrtf(1.0f); - ///scalblnf(0.0f, 1); + scalblnf(0.0f, 1); scalbnf(0.0f, 1); signbit(1.0f); sincosf(0.0f, &fX, &fY); - //sincospif(0.0f, &fX, &fY); + sincospif(0.0f, &fX, &fY); sinf(0.0f); sinhf(0.0f); sinpif(0.0f); @@ -115,9 +115,9 @@ __host__ void single_precision_math_functions() tanhf(0.0f); tgammaf(2.0f); truncf(0.0f); - ///y0f(1.0f); - ///y1f(1.0f); - ///ynf(1, 1.0f); + y0f(1.0f); + y1f(1.0f); + ynf(1, 1.0f); } static void compileOnHost() diff --git a/projects/hip/tests/src/experimental/xcompile/gapi.sh b/projects/hip/tests/src/experimental/xcompile/gapi.sh index 489171d65a..d1b3936dbf 100755 --- a/projects/hip/tests/src/experimental/xcompile/gapi.sh +++ b/projects/hip/tests/src/experimental/xcompile/gapi.sh @@ -1,9 +1,9 @@ #!/bin/bash -gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o +gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o -gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c hHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o hHipApi.o +gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c hHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o hHipApi.o -gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include hHipApi.o gHipApi.o ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -lm -o gApi +gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include hHipApi.o gHipApi.o ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -lm -o gApi diff --git a/projects/hip/tests/src/experimental/xcompile/ghipapi.sh b/projects/hip/tests/src/experimental/xcompile/ghipapi.sh index 611053896d..ac424871ef 100755 --- a/projects/hip/tests/src/experimental/xcompile/ghipapi.sh +++ b/projects/hip/tests/src/experimental/xcompile/ghipapi.sh @@ -1,8 +1,8 @@ #!/bin/bash -gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o +gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o -gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c hHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o hHipApi.o +gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c hHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o hHipApi.o hipcc hHipApi.o gHipApi.o -o gHipApi diff --git a/projects/hip/tests/src/experimental/xcompile/gxxapi.sh b/projects/hip/tests/src/experimental/xcompile/gxxapi.sh index 42569f93d2..9625f94380 100755 --- a/projects/hip/tests/src/experimental/xcompile/gxxapi.sh +++ b/projects/hip/tests/src/experimental/xcompile/gxxapi.sh @@ -1,8 +1,8 @@ #!/bin/bash -g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o +g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o -g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c hxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o hxxHipApi.o +g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c hxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o hxxHipApi.o -g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include gxxHipApi.o hxxHipApi.o ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxApi +g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include gxxHipApi.o hxxHipApi.o ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxApi diff --git a/projects/hip/tests/src/experimental/xcompile/gxxhipapi.sh b/projects/hip/tests/src/experimental/xcompile/gxxhipapi.sh index 2405b15cd1..690bccaacb 100755 --- a/projects/hip/tests/src/experimental/xcompile/gxxhipapi.sh +++ b/projects/hip/tests/src/experimental/xcompile/gxxhipapi.sh @@ -1,8 +1,8 @@ #!/bin/bash -g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o +g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o hipcc -c hxxHipApi.cpp -o hxxHipApi.o -g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include gxxHipApi.o hxxHipApi.o ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi +g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include gxxHipApi.o hxxHipApi.o ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi diff --git a/projects/hip/tests/src/experimental/xcompile/hipapig.sh b/projects/hip/tests/src/experimental/xcompile/hipapig.sh index c0f96d6180..47cf811bf9 100755 --- a/projects/hip/tests/src/experimental/xcompile/hipapig.sh +++ b/projects/hip/tests/src/experimental/xcompile/hipapig.sh @@ -1,3 +1,3 @@ #!/bin/bash -gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include gApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -lm -o gApi +gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include gApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -lm -o gApi diff --git a/projects/hip/tests/src/experimental/xcompile/hipapigxx.sh b/projects/hip/tests/src/experimental/xcompile/hipapigxx.sh index 832cc878a0..1b675964bd 100755 --- a/projects/hip/tests/src/experimental/xcompile/hipapigxx.sh +++ b/projects/hip/tests/src/experimental/xcompile/hipapigxx.sh @@ -1,3 +1,3 @@ #!/bin/bash -g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include gxxApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxApi +g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include gxxApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxApi diff --git a/projects/hip/tests/src/experimental/xcompile/hipg.sh b/projects/hip/tests/src/experimental/xcompile/hipg.sh index b0b6333874..799cc164b2 100755 --- a/projects/hip/tests/src/experimental/xcompile/hipg.sh +++ b/projects/hip/tests/src/experimental/xcompile/hipg.sh @@ -1,6 +1,6 @@ #!/bin/bash -gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o +gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o hipcc -c hHip.c -o hHip.o diff --git a/projects/hip/tests/src/experimental/xcompile/hipgapi.sh b/projects/hip/tests/src/experimental/xcompile/hipgapi.sh index 3804018719..d6919832e2 100755 --- a/projects/hip/tests/src/experimental/xcompile/hipgapi.sh +++ b/projects/hip/tests/src/experimental/xcompile/hipgapi.sh @@ -1,6 +1,6 @@ #!/bin/bash -gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o +gcc -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gHipApi.c ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gHipApi.o hipcc -c hHipApi.c -o hHipApi.o diff --git a/projects/hip/tests/src/experimental/xcompile/hipgxx.sh b/projects/hip/tests/src/experimental/xcompile/hipgxx.sh index 58b09c2ce9..19340cfcaf 100755 --- a/projects/hip/tests/src/experimental/xcompile/hipgxx.sh +++ b/projects/hip/tests/src/experimental/xcompile/hipgxx.sh @@ -1,6 +1,6 @@ #!/bin/bash -g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o +g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o hipcc -c hxxHip.cpp -o hxxHip.o diff --git a/projects/hip/tests/src/experimental/xcompile/hipgxxapi.sh b/projects/hip/tests/src/experimental/xcompile/hipgxxapi.sh index fd04425c47..8d3f61b941 100755 --- a/projects/hip/tests/src/experimental/xcompile/hipgxxapi.sh +++ b/projects/hip/tests/src/experimental/xcompile/hipgxxapi.sh @@ -1,6 +1,6 @@ #!/bin/bash -g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/staging_buffer.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o +g++ -D__HIP_PLATFORM_HCC__= -I${HIP_PATH}/include -I${HCC_HOME}/include -c gxxHipApi.cpp ${HIP_PATH}/lib/device_util.cpp.o ${HIP_PATH}/lib/hip_device.cpp.o ${HIP_PATH}/lib/hip_error.cpp.o ${HIP_PATH}/lib/hip_event.cpp.o ${HIP_PATH}/lib/hip_hcc.cpp.o ${HIP_PATH}/lib/hip_memory.cpp.o ${HIP_PATH}/lib/hip_peer.cpp.o ${HIP_PATH}/lib/hip_stream.cpp.o ${HIP_PATH}/lib/unpinned_copy_engine.cpp.o -L${HCC_HOME}/lib -lhc_am -L${HSA_PATH}/lib -lhsa-runtime64 -lc++ -lmcwamp -ldl -o gxxHipApi.o hipcc -c hxxHipApi.cpp -o hxxHipApi.o diff --git a/projects/hip/tests/src/hipHostAlloc.cpp b/projects/hip/tests/src/hipHostAlloc.cpp index 01ca04b311..a6c4cb20e0 100644 --- a/projects/hip/tests/src/hipHostAlloc.cpp +++ b/projects/hip/tests/src/hipHostAlloc.cpp @@ -1,24 +1,24 @@ /* -Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. + Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved. -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: + 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 above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + */ #include"test_common.h" @@ -26,43 +26,43 @@ THE SOFTWARE. #define SIZE LEN*sizeof(float) __global__ void Add(hipLaunchParm lp, float *Ad, float *Bd, float *Cd){ -int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; -Cd[tx] = Ad[tx] + Bd[tx]; + int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x; + Cd[tx] = Ad[tx] + Bd[tx]; } int main(){ -float *A, *B, *C; -float *Ad, *Bd, *Cd; + float *A, *B, *C; + float *Ad, *Bd, *Cd; -hipDeviceProp_t prop; -int device; -HIPCHECK(hipGetDevice(&device)); -HIPCHECK(hipGetDeviceProperties(&prop, device)); -if(prop.canMapHostMemory != 1){ -std::cout<<"Exiting..."< +#define GPU_PRINT_TIME(cmd, elapsed, quiet) do {\ + struct timeval start, stop;\ + float elapsed;\ + gettimeofday(&start, NULL);\ + hipDeviceSynchronize();\ + cmd;\ + hipDeviceSynchronize();\ + gettimeofday(&stop, NULL);\ + } while(0); + + + +#define MY_LAUNCH(command, doTrace, msg) \ +{\ + if (doTrace) printf ("TRACE: %s %s\n", msg, #command); \ + command;\ +} + + +#define MY_LAUNCH_WITH_PAREN(command, doTrace, msg) \ +{\ + if (doTrace) printf ("TRACE: %s %s\n", msg, #command); \ + (command);\ +} + + + int main() { float *Ad; @@ -32,5 +65,27 @@ int main() hipLaunchKernel(vAdd, 1024, dim3(1), 0, 0, Ad); hipLaunchKernel(vAdd, dim3(1024), 1, 0, 0, Ad); hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad); + + // Test case with hipLaunchKernel inside another macro: + float e0; + GPU_PRINT_TIME (hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), e0, j); + GPU_PRINT_TIME (WRAP(hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad)), e0, j); + +#ifdef EXTRA_PARENS_1 + // Don't wrap hipLaunchKernel in extra set of parens: + GPU_PRINT_TIME ((hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad)), e0, j); +#endif + + MY_LAUNCH (hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCall"); + + float *A; + float e1; + MY_LAUNCH_WITH_PAREN (hipMalloc(&A, 100), true, "launch2"); + +#ifdef EXTRA_PARENS_2 + //MY_LAUNCH_WITH_PAREN wraps cmd in () which can cause issues. + MY_LAUNCH_WITH_PAREN (hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCall"); +#endif + passed(); }