Merge branch 'amd-develop' into amd-master

Change-Id: Ie54323de830d1f8c4c12d3af03154b9fa5405ab3


[ROCm/hip commit: 74885ca71a]
This commit is contained in:
Maneesh Gupta
2016-08-19 16:16:45 +05:30
51 zmienionych plików z 3314 dodań i 1391 usunięć
+65 -23
Wyświetl plik
@@ -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)
+27
Wyświetl plik
@@ -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:
+43 -40
Wyświetl plik
@@ -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 (<CONFIG>) {
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 (<CONFIG>) {
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 ();
+40 -20
Wyświetl plik
@@ -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 (<CONFIG>) {
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");
}
+104 -15
Wyświetl plik
@@ -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<VarDecl>("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<BuiltinType>(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<VarDecl>("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 <hip_runtime.h>\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 <hip_runtime.h>\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);
@@ -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__".
@@ -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
@@ -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 <hip_runtime_api.h>
#include <hcblas.h>
//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
+179 -213
Wyświetl plik
@@ -22,7 +22,7 @@ THE SOFTWARE.
#include <hc.hpp>
#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<typename T>
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 <typename MUTEX_TYPE>
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 <typename MUTEX_TYPE>
class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE>
template <typename MUTEX_TYPE>
class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE>
{
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<ihipSignal_t> _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<StreamMutex> ihipStreamCritical_t;
typedef ihipStreamCriticalBase_t<StreamMutex> ihipStreamCritical_t;
typedef LockedAccessor<ihipStreamCritical_t> 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 MUTEX_TYPE>
class ihipDeviceCriticalBase_t : LockedBase<MUTEX_TYPE>
//----
// 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 <typename MUTEX_TYPE>
class ihipCtxCriticalBase_t : LockedBase<MUTEX_TYPE>
{
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<ihipDeviceCriticalBase_t>;
// Streams:
void addStream(ihipStream_t *stream);
std::list<ihipStream_t*> &streams() { return _streams; };
const std::list<ihipStream_t*> &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<ihipStream_t> > _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<ihipCtxCriticalBase_t>;
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<ihipDevice_t*> _peers; // list of enabled peer devices.
std::list<ihipCtx_t*> _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<DeviceMutex> ihipDeviceCritical_t;
// Note Mutex type Real/Fake selected based on CtxMutex
typedef ihipCtxCriticalBase_t<CtxMutex> ihipCtxCritical_t;
// This type is used by functions that need access to the critical device structures.
typedef LockedAccessor<ihipDeviceCritical_t> LockedAccessor_DeviceCrit_t;
typedef LockedAccessor<ihipCtxCritical_t> 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<ihipDeviceCritical_t> for examples; do not access _criticalData directly.
ihipDeviceCritical_t _criticalData;
// Search for LockedAccessor<ihipCtxCritical_t> 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<int> 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<typename T>
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<typename T>
hc::completion_future ihipMemsetKernel(hipStream_t, T*, T, size_t);
hipStream_t ihipSyncAndResolveStream(hipStream_t);
template <typename T>
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<sizeBytes; i+=stride) {
ptr[i] = val;
}
});
return cf;
os << "stream#";
os << s.getDevice()->_deviceId;;
os << '.';
os << s._id;
return os;
}
template <typename T>
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<sizeBytes; i+=stride) {
c[i] = a[i];
}
});
return cf;
os << '{';
os << s.x;
os << ',';
os << s.y;
os << ',';
os << s.z;
os << '}';
return os;
}
#endif
@@ -55,7 +55,9 @@ THE SOFTWARE.
#define USE_GRID_LAUNCH_20 0
#endif
#define HIP_LAUNCH_PARAM_BUFFER_POINTER ((void*) 0x01)
#define HIP_LAUNCH_PARAM_BUFFER_SIZE ((void*) 0x02)
#define HIP_LAUNCH_PARAM_END ((void*) 0x03)
extern int HIP_TRACE_API;
@@ -627,6 +629,8 @@ do {\
__attribute__((address_space(3))) type* var = \
(__attribute__((address_space(3))) type*)__get_dynamicgroupbaseptr(); \
#define HIP_DYNAMIC_SHARED_ATTRIBUTE __attribute__((address_space(3)))
#endif // __HCC__
@@ -43,7 +43,19 @@ THE SOFTWARE.
extern "C" {
#endif
//---
//API-visible structures
typedef struct ihipCtx_t *hipCtx_t;
// Note many APIs also use integer deviceIds as an alternative to the device pointer:
typedef struct ihipDevice_t *hipDevice_t;
typedef struct ihipStream_t *hipStream_t;
typedef uint64_t hipFunction;
typedef uint64_t hipModule;
typedef struct hipEvent_t {
struct ihipEvent_t *_handle;
} hipEvent_t;
@@ -203,7 +215,7 @@ hipError_t hipDeviceReset(void) ;
/**
* @brief Set default device to be used for subsequent hip API calls from this thread.
*
* @param[in] device Valid device in range 0...hipGetDeviceCount().
* @param[in] deviceId Valid device in range 0...hipGetDeviceCount().
*
* Sets @p device as the default device for the calling host thread. Valid device id's are 0... (hipGetDeviceCount()-1).
*
@@ -224,7 +236,7 @@ hipError_t hipDeviceReset(void) ;
*
* @see hipGetDevice, hipGetDeviceCount
*/
hipError_t hipSetDevice(int device);
hipError_t hipSetDevice(int deviceId);
/**
@@ -237,8 +249,10 @@ hipError_t hipSetDevice(int device);
* hipGetDevice returns in * @p device the default device for the calling host thread.
*
* @see hipSetDevice, hipGetDevicesizeBytes
*
* @returns hipSuccess, hipErrorInvalidDevice
*/
hipError_t hipGetDevice(int *device);
hipError_t hipGetDevice(int *deviceId);
/**
@@ -254,19 +268,19 @@ hipError_t hipGetDeviceCount(int *count);
* @brief Query device attribute.
* @param [out] pi pointer to value to return
* @param [in] attr attribute to query
* @param [in] device which device to query for information
* @param [in] deviceId which device to query for information
*/
hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device);
hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int deviceId);
/**
* @brief Returns device properties.
*
* @param [out] prop written with device properties
* @param [in] device which device to query for information
* @param [in] deviceId which device to query for information
*
* Populates hipGetDeviceProperties with information for the specified device.
*/
hipError_t hipGetDeviceProperties(hipDeviceProp_t* prop, int device);
hipError_t hipGetDeviceProperties(hipDeviceProp_t* prop, int deviceId);
@@ -378,14 +392,14 @@ const char *hipGetErrorName(hipError_t hip_error);
/**
* @brief Return handy text string message to explain the error which occurred
*
* @param hip_error Error code to convert to string.
* @param hipError Error code to convert to string.
* @return const char pointer to the NULL-terminated error string
*
* @warning : on HCC, this function returns the name of the error (same as hipGetErrorName)
*
* @see hipGetErrorName, hipGetLastError, hipPeakAtLastError, hipError_t
*/
const char *hipGetErrorString(hipError_t hip_error);
const char *hipGetErrorString(hipError_t hipError);
// end doxygen Error
/**
@@ -413,11 +427,10 @@ const char *hipGetErrorString(hipError_t hip_error);
* @return #hipSuccess, #hipErrorInvalidValue
*
* Create a new asynchronous stream. @p stream returns an opaque handle that can be used to reference the newly
* created stream in subsequent hipStream* commands. The stream is allocated on the heap and will remain allocated
* created stream in subsequent hipStream* commands. The stream is allocated on the heap and will remain allocated
*
* even if the handle goes out-of-scope. To release the memory used by the stream, applicaiton must call hipStreamDestroy.
* Flags controls behavior of the stream. See #hipStreamDefault, #hipStreamNonBlocking.
* @error hipStream_t are under development - with current HIP use the NULL stream.
*/
hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags);
@@ -431,12 +444,14 @@ hipError_t hipStreamCreateWithFlags(hipStream_t *stream, unsigned int flags);
* @return #hipSuccess, #hipErrorInvalidValue
*
* Create a new asynchronous stream. @p stream returns an opaque handle that can be used to reference the newly
* created stream in subsequent hipStream* commands. The stream is allocated on the heap and will remain allocated
* created stream in subsequent hipStream* commands. The stream is allocated on the heap and will remain allocated
* even if the handle goes out-of-scope. To release the memory used by the stream, applicaiton must call hipStreamDestroy.
*
*
*
* @see hipStreamDestroy
*
* @return
*
*/
hipError_t hipStreamCreate(hipStream_t *stream);
@@ -690,7 +705,7 @@ hipError_t hipMalloc(void** ptr, size_t size) ;
hipError_t hipMallocHost(void** ptr, size_t size) __attribute__((deprecated("use hipHostMalloc instead"))) ;
/**
* @brief Allocate device accessible page locked host memory
* @brief Allocate device accessible page locked host memory
*
* @param[out] ptr Pointer to the allocated host pinned memory
* @param[in] size Requested memory size
@@ -732,9 +747,9 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) ;
* - #hipHostRegisterMapped Map the allocation into the address space for the current device. The device pointer can be obtained with #hipHostGetDevicePointer.
*
*
* After registering the memory, use #hipHostGetDevicePointer to obtain the mapped device pointer.
* After registering the memory, use #hipHostGetDevicePointer to obtain the mapped device pointer.
* On many systems, the mapped device pointer will have a different value than the mapped host pointer. Applications
* must use the device pointer in device code, and the host pointer in device code.
* must use the device pointer in device code, and the host pointer in device code.
*
* On some systems, registered memory is pinned. On some systems, registered memory may not be actually be pinned
* but uses OS or hardware facilities to all GPU access to the host memory.
@@ -742,7 +757,7 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) ;
* Developers are strongly encouraged to register memory blocks which are aligned to the host cache-line size.
* (typically 64-bytes but can be obtains from the CPUID instruction).
*
* If registering non-aligned pointers, the application must take care when register pointers from the same cache line
* If registering non-aligned pointers, the application must take care when register pointers from the same cache line
* on different devices. HIP's coarse-grained synchronization model does not guarantee correct results if different
* devices write to different parts of the same cache block - typically one of the writes will "win" and overwrite data
* from the other registered memory region.
@@ -780,7 +795,7 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height
* If pointer is NULL, the hip runtime is initialized and hipSuccess is returned.
*
* @param[in] ptr Pointer to memory to be freed
* @return #hipSuccess
* @return #hipSuccess
* @return #hipErrorInvalidDevicePointer (if pointer is invalid, including host pointers allocated with hipHostMalloc)
*/
hipError_t hipFree(void* ptr);
@@ -801,7 +816,7 @@ hipError_t hipFreeHost(void* ptr) __attribute__((deprecated("use hipHostFree ins
* If pointer is NULL, the hip runtime is initialized and hipSuccess is returned.
*
* @param[in] ptr Pointer to memory to be freed
* @return #hipSuccess,
* @return #hipSuccess,
* #hipErrorInvalidValue (if pointer is invalid, including device pointers allocated with hipMalloc)
*/
hipError_t hipHostFree(void* ptr);
@@ -817,7 +832,7 @@ hipError_t hipHostFree(void* ptr);
*
* For hipMemcpy, the copy is always performed by the current device (set by hipSetDevice).
* For multi-gpu or peer-to-peer configurations, it is recommended to set the current device to the device where the src data is physically located.
* For optimal peer-to-peer copies, the copy device must be able to access the src and dst pointers (by calling hipDeviceEnablePeerAccess with copy agent as the
* For optimal peer-to-peer copies, the copy device must be able to access the src and dst pointers (by calling hipDeviceEnablePeerAccess with copy agent as the
* current device and src/dest as the peerDevice argument. if this is not done, the hipMemcpy will still work, but will perform the copy using a staging buffer
* on the host.
*
@@ -835,7 +850,7 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
*
* The memory areas may not overlap. Symbol can either be a variable that resides in global or constant memory space, or it can be a character string,
* naming a variable that resides in global or constant memory space. Kind can be either hipMemcpyHostToDevice or hipMemcpyDeviceToDevice
* TODO: cudaErrorInvalidSymbol and cudaErrorInvalidMemcpyDirection is not supported, use hipErrorUnknown for now.
* TODO: cudaErrorInvalidSymbol and cudaErrorInvalidMemcpyDirection is not supported, use hipErrorUnknown for now.
*
* @param[in] symbolName - Symbol destination on device
* @param[in] src - Data being copy from
@@ -856,7 +871,7 @@ hipError_t hipMemcpyToSymbol(const char* symbolName, const void *src, size_t siz
* For hipMemcpy, the copy is always performed by the device associated with the specified stream.
*
* For multi-gpu or peer-to-peer configurations, it is recommended to use a stream which is a attached to the device where the src data is physically located.
* For optimal peer-to-peer copies, the copy device must be able to access the src and dst pointers (by calling hipDeviceEnablePeerAccess with copy agent as the
* For optimal peer-to-peer copies, the copy device must be able to access the src and dst pointers (by calling hipDeviceEnablePeerAccess with copy agent as the
* current device and src/dest as the peerDevice argument. if this is not done, the hipMemcpy will still work, but will perform the copy using a staging buffer
* on the host.
*
@@ -943,7 +958,7 @@ hipError_t hipMemGetInfo (size_t * free, size_t * total) ;
*
* Returns "0" in @p canAccessPeer if deviceId == peerDeviceId, and both are valid devices : a device is not a peer of itself.
*
* @returns #hipSuccess,
* @returns #hipSuccess,
* @returns #hipErrorInvalidDevice if deviceId or peerDeviceId are not valid devices
* @warning PeerToPeer support is experimental.
*/
@@ -951,7 +966,7 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDev
/**
* @brief Enable direct access from current device's virtual address space to memory allocations physically located on a peer device.
* @brief Enable direct access from current device's virtual address space to memory allocations physically located on a peer device.
*
* Memory which already allocated on peer device will be mapped into the address space of the current device. In addition, all
* future memory allocations on peerDeviceId will be mapped into the address space of the current device when the memory is allocated.
@@ -961,7 +976,7 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDev
* @param [in] peerDeviceId
* @param [in] flags
*
* Returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue,
* Returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue,
* @returns #hipErrorPeerAccessAlreadyEnabled if peer access is already enabled for this device.
* @warning PeerToPeer support is experimental.
*/
@@ -969,7 +984,7 @@ hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags);
/**
* @brief Disable direct access from current device's virtual address space to memory allocations physically located on a peer device.
* @brief Disable direct access from current device's virtual address space to memory allocations physically located on a peer device.
*
* Returns hipErrorPeerAccessNotEnabled if direct access to memory on peerDevice has not yet been enabled from the current device.
*
@@ -1021,16 +1036,57 @@ hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int src
* @}
*/
/**
*-------------------------------------------------------------------------------------------------
*-------------------------------------------------------------------------------------------------
* @defgroup Version Management
* @defgroup Driver Initialization and Version
* @{
*
*/
/**
* @brief Explicitly initializes the HIP runtime.
*
* Most HIP APIs implicitly initialize the HIP runtime.
* This API provides control over the timing of the initialization.
*/
// TODO-ctx - more description on error codes.
hipError_t hipInit(unsigned int flags) ;
// TODO-ctx
hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device);
hipError_t hipCtxDestroy(hipCtx_t ctx);
hipError_t hipCtxPopCurrent(hipCtx_t* ctx);
hipError_t hipCtxPushCurrent(hipCtx_t ctx);
hipError_t hipCtxSetCurrent(hipCtx_t ctx);
hipError_t hipCtxGetCurrent(hipCtx_t* ctx);
hipError_t hipCtxGetDevice(hipDevice_t *device);
hipError_t hipCtxGetApiVersion (hipCtx_t ctx,int *apiVersion);
hipError_t hipCtxGetCacheConfig ( hipFuncCache *cacheConfig );
hipError_t hipCtxSetCacheConfig ( hipFuncCache cacheConfig );
hipError_t hipCtxSetSharedMemConfig ( hipSharedMemConfig config );
hipError_t hipCtxGetSharedMemConfig ( hipSharedMemConfig * pConfig );
// TODO-ctx
/**
* @return hipSuccess, hipErrorInvalidDevice
*/
hipError_t hipDeviceGetFromId(hipDevice_t *device, int deviceId);
/**
* @brief Returns the approximate HIP driver version.
*
@@ -1044,6 +1100,21 @@ hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int src
hipError_t hipDriverGetVersion(int *driverVersion) ;
hipError_t hipModuleLoad(hipModule *module, const char *fname);
hipError_t hipModuleGetFunction(hipFunction *function, hipModule module, const char *kname);
hipError_t hipLaunchModuleKernel(hipFunction f,
unsigned int gridDimX,
unsigned int gridDimY,
unsigned int gridDimZ,
unsigned int blockDimX,
unsigned int blockDimY,
unsigned int blockDimZ,
unsigned int sharedMemBytes,
hipStream_t stream,
void **kernelParams,
void **extra) __attribute__((deprecated("kernelParams is not fully supported, use extra instead"))) ;
// doxygen end Version Management
/**
@@ -26,42 +26,46 @@ THE SOFTWARE.
//-------------------------------------------------------------------------------------------------
// An optimized "staging buffer" used to implement Host-To-Device and Device-To-Host copies.
// Some GPUs may not be able to directly access host memory, and in these cases we need to
// Some GPUs may not be able to directly access host memory, and in these cases we need to
// stage the copy through a pinned staging buffer. For example, the CopyHostToDevice
// uses the CPU to copy to a pinned "staging buffer", and then use the GPU DMA engine to copy
// from the staging buffer to the final destination. The copy is broken into buffer-sized chunks
// to limit the size of the buffer and also to provide better performance by overlapping the CPU copies
// to limit the size of the buffer and also to provide better performance by overlapping the CPU copies
// with the DMA copies.
//
// PinInPlace is another algorithm which pins the host memory "in-place", and copies it with the DMA
// engine. This routine is under development.
//
// Staging buffer provides thread-safe access via a mutex.
struct StagingBuffer {
struct UnpinnedCopyEngine {
static const int _max_buffers = 4;
StagingBuffer(hsa_agent_t hsaAgent, hsa_region_t systemRegion, size_t bufferSize, int numBuffers) ;
~StagingBuffer();
UnpinnedCopyEngine(hsa_agent_t hsaAgent,hsa_agent_t cpuAgent, size_t bufferSize, int numBuffers,int thresholdH2D_directStaging,int thresholdH2D_stagingPinInPlace,int thresholdD2H) ;
~UnpinnedCopyEngine();
void CopyHostToDevice(void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor);
void CopyHostToDevice(int tempIndex,int isLargeBar,void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor);
void CopyHostToDevicePinInPlace(void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor);
void CopyDeviceToHost (void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor);
void CopyDeviceToHost (int tempIndex,void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor);
void CopyDeviceToHostPinInPlace(void* dst, const void* src, size_t sizeBytes, hsa_signal_t *waitFor);
void CopyPeerToPeer( void* dst, hsa_agent_t dstAgent, const void* src, hsa_agent_t srcAgent, size_t sizeBytes, hsa_signal_t *waitFor);
private:
hsa_agent_t _hsa_agent;
hsa_agent_t _hsaAgent;
hsa_agent_t _cpuAgent;
size_t _bufferSize; // Size of the buffers.
int _numBuffers;
char *_pinnedStagingBuffer[_max_buffers];
hsa_signal_t _completion_signal[_max_buffers];
hsa_signal_t _completion_signal2[_max_buffers]; // P2P needs another set of signals.
std::mutex _copy_lock; // provide thread-safe access
std::mutex _copy_lock; // provide thread-safe access
int _hipH2DTransferThresholdDirectOrStaging;
int _hipH2DTransferThresholdStagingOrPininplace;
int _hipD2HTransferThreshold;
};
#endif
+67 -19
Wyświetl plik
@@ -32,6 +32,13 @@ THE SOFTWARE.
#include <string.h> // for getDeviceProp
#include <hip/hip_common.h>
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;
/*
+66
Wyświetl plik
@@ -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 <hcc_detail/hip_blas.h>
#elif defined(__HIP_PLATFORM_NVCC__) and not defined (__HIP_PLATFORM_HCC__)
#include <nvcc_detail/hip_blas.h>
#else
#error("Must define exactly one of __HIP_PLATFORM_HCC__ or __HIP_PLATFORM_NVCC__");
#endif
@@ -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 <cuda_runtime_api.h>
#include <cublas.h>
#include <cublas_v2.h>
//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
@@ -99,6 +99,8 @@ kernelName<<<numblocks,numthreads,memperblock,streamId>>>(0, ##__VA_ARGS__);\
#define HIP_DYNAMIC_SHARED(type, var) \
extern __shared__ type var[]; \
#define HIP_DYNAMIC_SHARED_ATTRIBUTE
#endif
@@ -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<class T>
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;
+1
Wyświetl plik
@@ -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
@@ -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
@@ -0,0 +1,94 @@
#include <random>
#include <algorithm>
#include <iostream>
#include <cmath>
// header file for the GPU API
#include <cuda_runtime.h>
#include <cublas_v2.h>
#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<float> 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;
}
@@ -0,0 +1,94 @@
#include <random>
#include <algorithm>
#include <iostream>
#include <cmath>
// header file for the GPU API
#include <hip_runtime.h>
#include <hipblas.h>
#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<float> 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;
}
+1 -1
Wyświetl plik
@@ -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;
}
+219
Wyświetl plik
@@ -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 <stack>
#include "hip_runtime.h"
#include "hcc_detail/hip_hcc.h"
#include "hcc_detail/trace_helper.h"
// Stack of contexts
thread_local std::stack<ihipCtx_t *> 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);
}
+54 -22
Wyświetl plik
@@ -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);
}
+2 -3
Wyświetl plik
@@ -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);
}
+7 -7
Wyświetl plik
@@ -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);
}
Plik diff jest za duży Load Diff
+307 -223
Wyświetl plik
@@ -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::vector<hc::accelerator>vecAcc;
for(int i=0;i<g_deviceCnt;i++){
vecAcc.push_back(g_devices[i]._acc);
vecAcc.push_back(ihipGetDevice(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 <typename T>
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<sizeBytes; i+=stride) {
ptr[i] = val;
}
});
return cf;
}
template <typename T>
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; i<sizeBytes; i+=stride) {
c[i] = a[i];
}
});
return cf;
}
// TODO-sync: function is async unless target is pinned host memory - then these are fully sync.
/** @return #hipErrorInvalidValue
*/
*/
hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream )
{
HIP_INIT_API(dst, value, sizeBytes, stream);
@@ -692,19 +775,20 @@ hipError_t hipMemGetInfo (size_t *free, size_t *total)
hipError_t e = hipSuccess;
ihipDevice_t * hipDevice = ihipGetTlsDefaultDevice();
if (hipDevice) {
ihipCtx_t * ctx = ihipGetTlsDefaultCtx();
if (ctx) {
auto device = ctx->getWriteableDevice();
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
+254
Wyświetl plik
@@ -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 <fstream>
//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"<<std::endl;
}
hsa_code_object_t obj;
in.seekg(0, std::ios::beg);
std::copy(std::istreambuf_iterator<char>(in),
std::istreambuf_iterator<char>(), 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;
}
+76 -45
Wyświetl plik
@@ -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);
}
+25 -19
Wyświetl plik
@@ -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;
@@ -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<std::mutex> 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<char*> (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<char*> (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<char*> (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<std::mutex> l (_copy_lock);
const char *srcp = static_cast<const char*> (src);
char *dstp = static_cast<char*> (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<std::mutex> 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<const char*> (src);
char *dstp = static_cast<char*> (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<std::mutex> 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<char*> (dstp), theseBytes, &_hsa_agent, 1, &locked_destp);
hsa_status_t hsa_status = hsa_amd_memory_lock(const_cast<char*> (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<std::mutex> l (_copy_lock);
const char *srcp0 = static_cast<const char*> (src);
char *dstp1 = static_cast<char*> (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<std::mutex> l (_copy_lock);
if (sizeBytes >= UINT64_MAX/2) {
THROW_ERROR (hipErrorInvalidValue);
}
const char *srcp0 = static_cast<const char*> (src);
char *dstp1 = static_cast<char*> (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<std::mutex> 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;
+5 -2
Wyświetl plik
@@ -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)
@@ -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 " " )
@@ -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();
};
@@ -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)
@@ -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()
@@ -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)
@@ -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()
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
@@ -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
+44 -44
Wyświetl plik
@@ -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..."<<std::endl;
failed("Does support HostPinned Memory");
}
hipDeviceProp_t prop;
int device;
HIPCHECK(hipGetDevice(&device));
HIPCHECK(hipGetDeviceProperties(&prop, device));
if(prop.canMapHostMemory != 1){
std::cout<<"Exiting..."<<std::endl;
failed("Does support HostPinned Memory");
}
HIPCHECK(hipHostMalloc((void**)&A, SIZE, hipHostMallocWriteCombined | hipHostMallocMapped));
HIPCHECK(hipHostMalloc((void**)&B, SIZE, hipHostMallocDefault));
HIPCHECK(hipHostMalloc((void**)&C, SIZE, hipHostMallocMapped));
HIPCHECK(hipHostMalloc((void**)&A, SIZE, hipHostMallocWriteCombined | hipHostMallocMapped));
HIPCHECK(hipHostMalloc((void**)&B, SIZE, hipHostMallocDefault));
HIPCHECK(hipHostMalloc((void**)&C, SIZE, hipHostMallocMapped));
HIPCHECK(hipHostGetDevicePointer((void**)&Ad, A, 0));
HIPCHECK(hipHostGetDevicePointer((void**)&Cd, C, 0));
HIPCHECK(hipHostGetDevicePointer((void**)&Ad, A, 0));
HIPCHECK(hipHostGetDevicePointer((void**)&Cd, C, 0));
for(int i=0;i<LEN;i++){
A[i] = 1.0f;
B[i] = 2.0f;
}
for(int i=0;i<LEN;i++){
A[i] = 1.0f;
B[i] = 2.0f;
}
HIPCHECK(hipMalloc((void**)&Bd, SIZE));
HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
HIPCHECK(hipMalloc((void**)&Bd, SIZE));
HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
dim3 dimGrid(LEN/512,1,1);
dim3 dimBlock(512,1,1);
dim3 dimGrid(LEN/512,1,1);
dim3 dimBlock(512,1,1);
hipLaunchKernel(HIP_KERNEL_NAME(Add), dimGrid, dimBlock, 0, 0, Ad, Bd, Cd);
hipLaunchKernel(HIP_KERNEL_NAME(Add), dimGrid, dimBlock, 0, 0, Ad, Bd, Cd);
passed();
passed();
}
@@ -24,6 +24,39 @@ THE SOFTWARE.
__global__ void vAdd(hipLaunchParm lp, float *a){}
//---
//Some wrapper macro for testing:
#define WRAP(...) __VA_ARGS__
#include <sys/time.h>
#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();
}