Merge 'master' into 'amd-master'
Change-Id: Id4dae3f5d145db304d2bd60a1105dd9563805da9
[ROCm/hip commit: 2a1542ab9d]
This commit is contained in:
Jenkins
Vendored
+4
-4
@@ -295,13 +295,13 @@ def docker_upload_dockerhub( String local_org, String image_name, String remote_
|
||||
String build_config = 'Release'
|
||||
String job_name = env.JOB_NAME.toLowerCase( )
|
||||
|
||||
// The following launches 3 builds in parallel: rocm-head, rocm-2.7.x and cuda-10.x
|
||||
parallel rocm_2_7:
|
||||
// The following launches 3 builds in parallel: rocm-head, rocm-2.8.x and cuda-10.x
|
||||
parallel rocm_2_8:
|
||||
{
|
||||
node('hip-rocm')
|
||||
{
|
||||
String hcc_ver = 'rocm-2.7.x'
|
||||
String from_image = 'ci_test_nodes/rocm-2.7.x/ubuntu-16.04:latest'
|
||||
String hcc_ver = 'rocm-2.8.x'
|
||||
String from_image = 'ci_test_nodes/rocm-2.8.x/ubuntu-16.04:latest'
|
||||
String inside_args = '--device=/dev/kfd --device=/dev/dri --group-add=video'
|
||||
|
||||
// Checkout source code, dependencies and version files
|
||||
|
||||
@@ -1,7 +1,7 @@
|
||||
#!/usr/bin/perl -w
|
||||
|
||||
$HIP_BASE_VERSION_MAJOR = "2";
|
||||
$HIP_BASE_VERSION_MINOR = "8";
|
||||
$HIP_BASE_VERSION_MINOR = "10";
|
||||
|
||||
# Need perl > 5.10 to use logic-defined or
|
||||
use 5.006; use v5.10.1;
|
||||
|
||||
+1701
-1699
File diff suppressed because it is too large
Load Diff
@@ -599,9 +599,21 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DEVICE_FUNC_MAP{
|
||||
{"h2exp10", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"h2cos", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"h2sin", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"__shfl_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__shfl_up_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__shfl_down_sync",{"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__shfl_xor_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"atomicAdd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__shfl_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__shfl_up_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__shfl_down_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
{"__shfl_xor_sync", {"", "", CONV_DEVICE_FUNC, API_RUNTIME, UNSUPPORTED}},
|
||||
// atomic functions
|
||||
{"atomicAdd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"atomicSub", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"atomicExch", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"atomicMin", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"atomicMax", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"atomicInc", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"atomicDec", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"atomicDec", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"atomicAnd", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"atomicOr", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"atomicXor", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
{"atomicCAS", {"", "", CONV_DEVICE_FUNC, API_RUNTIME}},
|
||||
};
|
||||
|
||||
@@ -64,9 +64,11 @@ namespace perl {
|
||||
const std::string triple_tab = double_tab + tab;
|
||||
const std::string sSub = "sub";
|
||||
const std::string sReturn_0 = "return 0;\n";
|
||||
const std::string sReturn_m = "return $m;\n";
|
||||
const std::string sReturn_k = "return $k;\n";
|
||||
const std::string sForeach = "foreach $func (\n";
|
||||
const std::string sMy = "my $m = 0;\n";
|
||||
const std::string sMy_k = "my $k = 0;";
|
||||
const std::string sNoWarns = "no warnings qw/uninitialized/;";
|
||||
|
||||
const std::string sCudaDevice = "cudaDevice";
|
||||
const std::string sCudaDeviceId = "cudaDeviceId";
|
||||
const std::string sCudaDevices = "cudaDevices";
|
||||
@@ -142,21 +144,98 @@ namespace perl {
|
||||
*streamPtr.get() << tab << "foreach $stat (@statNames) {" << std::endl;
|
||||
*streamPtr.get() << double_tab << "printf STDERR \"%s:%d \", $stat, $counts{$stat};" << std::endl;
|
||||
*streamPtr.get() << tab << "}" << std::endl;
|
||||
*streamPtr.get() << tab << "printf STDERR \")\\n warn:%d LOC:%d\", $warnings, $loc;" << std::endl << "};" << std::endl;
|
||||
*streamPtr.get() << tab << "printf STDERR \")\\n warn:%d LOC:%d\", $warnings, $loc;" << std::endl << "}" << std::endl;
|
||||
for (int i = 0; i < 2; ++i) {
|
||||
*streamPtr.get() << std::endl << sSub << " " << (i ? "clearStats" : "addStats") << " {" << std::endl;
|
||||
*streamPtr.get() << tab << "my $dest_ref = shift();" << std::endl;
|
||||
*streamPtr.get() << tab << (i ? "my @statNames = @{ shift() };" : "my %adder = %{ shift() };") << std::endl;
|
||||
*streamPtr.get() << tab << "foreach " << (i ? "$stat(@statNames)" : "$key (keys %adder)") << " {" << std::endl;
|
||||
*streamPtr.get() << double_tab << "$dest_ref->" << (i ? "{$stat} = 0;" : "{$key} += $adder{$key};") << std::endl << tab << "}" << std::endl << "};" << std::endl;
|
||||
*streamPtr.get() << double_tab << "$dest_ref->" << (i ? "{$stat} = 0;" : "{$key} += $adder{$key};") << std::endl << tab << "}" << std::endl << "}" << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
void generateSimpleSubstitutions(std::unique_ptr<std::ostream>& streamPtr) {
|
||||
*streamPtr.get() << std::endl << sSub << " simpleSubstitutions" << " {" << std::endl;
|
||||
for (int i = 0; i < NUM_CONV_TYPES; ++i) {
|
||||
if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) {
|
||||
for (auto& ma : CUDA_INCLUDE_MAP) {
|
||||
if (Statistics::isUnsupported(ma.second)) continue;
|
||||
if (i == ma.second.type) {
|
||||
std::string sCUDA = ma.first.str();
|
||||
std::string sHIP = ma.second.hipName.str();
|
||||
sCUDA = std::regex_replace(sCUDA, std::regex("/"), "\\/");
|
||||
sHIP = std::regex_replace(sHIP, std::regex("/"), "\\/");
|
||||
*streamPtr.get() << tab << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << sCUDA << "\\b/" << sHIP << "/g;" << std::endl;
|
||||
}
|
||||
}
|
||||
} else {
|
||||
for (auto& ma : CUDA_RENAMES_MAP()) {
|
||||
if (Statistics::isUnsupported(ma.second)) continue;
|
||||
if (i == ma.second.type) {
|
||||
*streamPtr.get() << tab << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << ma.first.str() << "\\b/" << ma.second.hipName.str() << "/g;" << std::endl;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
*streamPtr.get() << "}" << std::endl;
|
||||
}
|
||||
|
||||
void generateExternShared(std::unique_ptr<std::ostream>& streamPtr) {
|
||||
*streamPtr.get() << std::endl << "# CUDA extern __shared__ syntax replace with HIP_DYNAMIC_SHARED() macro" << std::endl;
|
||||
*streamPtr.get() << sSub << " transformExternShared" << " {" << std::endl;
|
||||
*streamPtr.get() << tab << sNoWarns << std::endl;
|
||||
*streamPtr.get() << tab << sMy_k << std::endl;
|
||||
*streamPtr.get() << tab << "$k += s/extern\\s+([\\w\\(\\)]+)?\\s*__shared__\\s+([\\w:<>\\s]+)\\s+(\\w+)\\s*\\[\\s*\\]\\s*;/HIP_DYNAMIC_SHARED($1 $2, $3)/g;" << std::endl;
|
||||
*streamPtr.get() << tab << "$ft{ 'extern_shared' } += $k;" << std::endl << "}" << std::endl;
|
||||
}
|
||||
|
||||
void generateKernelLaunch(std::unique_ptr<std::ostream>& streamPtr) {
|
||||
*streamPtr.get() << std::endl << "# CUDA Kernel Launch Syntax" << std::endl;
|
||||
*streamPtr.get() << sSub << " transformKernelLaunch" << " {" << std::endl;
|
||||
*streamPtr.get() << tab << "my $TkernRef = @_;" << std::endl;
|
||||
*streamPtr.get() << tab << sNoWarns << std::endl;
|
||||
*streamPtr.get() << tab << sMy_k << std::endl << std::endl;
|
||||
|
||||
*streamPtr.get() << tab << "# Handle the kern<...><<<Dg, Db, Ns, S>>>() syntax with empty args:" << std::endl;
|
||||
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, $6)/g;" << std::endl;
|
||||
*streamPtr.get() << tab << "# Handle the kern<<<Dg, Db, Ns, S>>>() syntax with empty args:" << std::endl;
|
||||
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, $5)/g;" << std::endl << std::endl;
|
||||
|
||||
*streamPtr.get() << tab << "# Handle the kern<...><<<Dg, Db, Ns, S>>>(...) syntax with non-empty args:" << std::endl;
|
||||
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, $6, /g;" << std::endl;
|
||||
*streamPtr.get() << tab << "# Handle the kern<<<Dg, Db, Ns, S>>>(...) syntax with non-empty args:" << std::endl;
|
||||
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, $5, /g;" << std::endl << std::endl;
|
||||
|
||||
*streamPtr.get() << tab << "# Handle the kern<...><<<Dg, Db, Ns>>>() syntax with empty args:" << std::endl;
|
||||
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, 0)/g;" << std::endl;
|
||||
*streamPtr.get() << tab << "# Handle the kern<<<Dg, Db, Ns>>>() syntax with empty args:" << std::endl;
|
||||
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, 0)/g;" << std::endl << std::endl;
|
||||
|
||||
*streamPtr.get() << tab << "# Handle the kern<...><<Dg, Db, Ns>>>(...) syntax with non-empty args:" << std::endl;
|
||||
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, 0, /g;" << std::endl;
|
||||
*streamPtr.get() << tab << "# Handle the kern<<<Dg, Db, Ns>>>(...) syntax with non-empty args:" << std::endl;
|
||||
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, 0, /g;" << std::endl << std::endl;
|
||||
|
||||
*streamPtr.get() << tab << "# Handle the kern<...><<<Dg, Db>>>() syntax with empty args:" << std::endl;
|
||||
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), 0, 0)/g;" << std::endl;
|
||||
*streamPtr.get() << tab << "# Handle the kern<<<Dg, Db>>>() syntax with empty args:" << std::endl;
|
||||
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0)/g;" << std::endl << std::endl;
|
||||
|
||||
*streamPtr.get() << tab << "# Handle the kern<...><<<Dg, Db>>>(...) syntax with non-empty args:" << std::endl;
|
||||
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), 0, 0, /g;" << std::endl;
|
||||
*streamPtr.get() << tab << "# Handle the kern<<<Dg, Db>>>(...) syntax with non-empty args:" << std::endl;
|
||||
*streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0, /g;" << std::endl << std::endl;
|
||||
|
||||
*streamPtr.get() << tab << "if ($k) {" << std::endl;
|
||||
*streamPtr.get() << double_tab << "$ft{'kernel_launch'} += $k;" << std::endl;
|
||||
*streamPtr.get() << double_tab << "@$TkernRef{$1} ++;" << std::endl << tab << "}" << std::endl << "}" << std::endl;
|
||||
}
|
||||
|
||||
void generateHostFunctions(std::unique_ptr<std::ostream>& streamPtr) {
|
||||
*streamPtr.get() << std::endl << sSub << " transformHostFunctions" << "{" << std::endl << tab << sMy;
|
||||
*streamPtr.get() << std::endl << sSub << " transformHostFunctions" << " {" << std::endl << tab << sMy_k << std::endl;
|
||||
std::set<std::string> &funcSet = DeviceSymbolFunctions0;
|
||||
const std::string s0 = "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,/$func\\(";
|
||||
const std::string s1 = "$m += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,\\s*([^,\\)]+)\\s*(,\\s*|\\))\\s*/$func\\($2, ";
|
||||
const std::string s0 = "$k += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,/$func\\(";
|
||||
const std::string s1 = "$k += s/(?<!\\/\\/ CHECK: )($func)\\s*\\(\\s*([^,]+)\\s*,\\s*([^,\\)]+)\\s*(,\\s*|\\))\\s*/$func\\($2, ";
|
||||
for (int i = 0; i < 4; ++i) {
|
||||
*streamPtr.get() << tab + sForeach;
|
||||
switch (i) {
|
||||
@@ -187,7 +266,7 @@ namespace perl {
|
||||
}
|
||||
*streamPtr.get() << tab << "}" << std::endl;
|
||||
}
|
||||
*streamPtr.get() << tab << sReturn_m << "}" << std::endl;
|
||||
*streamPtr.get() << tab << sReturn_k << "}" << std::endl;
|
||||
}
|
||||
|
||||
void generateDeviceFunctions(std::unique_ptr<std::ostream>& streamPtr) {
|
||||
@@ -204,7 +283,7 @@ namespace perl {
|
||||
std::stringstream subCountSupported;
|
||||
std::stringstream subWarnUnsupported;
|
||||
std::stringstream subCommon;
|
||||
std::string sCommon = tab + sMy + tab + sForeach;
|
||||
std::string sCommon = tab + sMy_k + "\n" + tab + sForeach;
|
||||
subCountSupported << std::endl << sSub << " countSupportedDeviceFunctions" << " {" << std::endl << (countSupported ? sCommon : tab + sReturn_0);
|
||||
subWarnUnsupported << std::endl << sSub << " warnUnsupportedDeviceFunctions" << " {" << std::endl << (countUnsupported ? tab + "my $line_num = shift;\n" + sCommon : tab + sReturn_0);
|
||||
if (countSupported) {
|
||||
@@ -220,7 +299,7 @@ namespace perl {
|
||||
subCommon << double_tab << "my $mt_namespace = m/(\\w+)::($func)\\s*\\(\\s*.*\\s*\\)/g;" << std::endl;
|
||||
subCommon << double_tab << "my $mt = m/($func)\\s*\\(\\s*.*\\s*\\)/g;" << std::endl;
|
||||
subCommon << double_tab << "if ($mt && !$mt_namespace) {" << std::endl;
|
||||
subCommon << triple_tab << "$m += $mt;" << std::endl;
|
||||
subCommon << triple_tab << "$k += $mt;" << std::endl;
|
||||
}
|
||||
if (countSupported) {
|
||||
subCountSupported << subCommon.str();
|
||||
@@ -230,7 +309,7 @@ namespace perl {
|
||||
subWarnUnsupported << triple_tab << "print STDERR \" warning: $fileName:$line_num: unsupported device function \\\"$func\\\": $_\\n\";" << std::endl;
|
||||
}
|
||||
if (countSupported || countUnsupported) {
|
||||
sCommon = double_tab + "}\n" + tab + "}\n" + tab + sReturn_m;
|
||||
sCommon = double_tab + "}\n" + tab + "}\n" + tab + sReturn_k;
|
||||
}
|
||||
if (countSupported) subCountSupported << sCommon;
|
||||
if (countUnsupported) subWarnUnsupported << sCommon;
|
||||
@@ -267,29 +346,9 @@ namespace perl {
|
||||
}
|
||||
*streamPtr.get() << "\"" << counterNames[NUM_CONV_TYPES - 1] << "\");" << std::endl;
|
||||
generateStatFunctions(streamPtr);
|
||||
*streamPtr.get() << std::endl << sConv << "$ft{'" << counterNames[NUM_CONV_TYPES - 1] << "'};" << std::endl << std::endl;
|
||||
for (int i = 0; i < NUM_CONV_TYPES; ++i) {
|
||||
if (i == CONV_INCLUDE_CUDA_MAIN_H || i == CONV_INCLUDE) {
|
||||
for (auto& ma : CUDA_INCLUDE_MAP) {
|
||||
if (Statistics::isUnsupported(ma.second)) continue;
|
||||
if (i == ma.second.type) {
|
||||
std::string sCUDA = ma.first.str();
|
||||
std::string sHIP = ma.second.hipName.str();
|
||||
sCUDA = std::regex_replace(sCUDA, std::regex("/"), "\\/");
|
||||
sHIP = std::regex_replace(sHIP, std::regex("/"), "\\/");
|
||||
*streamPtr.get() << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << sCUDA << "\\b/" << sHIP << "/g;" << std::endl;
|
||||
}
|
||||
}
|
||||
}
|
||||
else {
|
||||
for (auto& ma : CUDA_RENAMES_MAP()) {
|
||||
if (Statistics::isUnsupported(ma.second)) continue;
|
||||
if (i == ma.second.type) {
|
||||
*streamPtr.get() << "$ft{'" << counterNames[ma.second.type] << "'} += s/\\b" << ma.first.str() << "\\b/" << ma.second.hipName.str() << "/g;" << std::endl;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
generateSimpleSubstitutions(streamPtr);
|
||||
generateExternShared(streamPtr);
|
||||
generateKernelLaunch(streamPtr);
|
||||
generateHostFunctions(streamPtr);
|
||||
generateDeviceFunctions(streamPtr);
|
||||
streamPtr.get()->flush();
|
||||
|
||||
@@ -38,6 +38,7 @@ namespace mat = clang::ast_matchers;
|
||||
|
||||
const std::string sHIP_DYNAMIC_SHARED = "HIP_DYNAMIC_SHARED";
|
||||
std::string sHIP_SYMBOL = "HIP_SYMBOL";
|
||||
std::string sHIP_KERNEL_NAME = "HIP_KERNEL_NAME";
|
||||
std::string s_reinterpret_cast = "reinterpret_cast<const void*>";
|
||||
const std::string sHipLaunchKernelGGL = "hipLaunchKernelGGL(";
|
||||
const std::string sDim3 = "dim3(";
|
||||
@@ -348,7 +349,7 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc
|
||||
clang::LangOptions DefaultLangOptions;
|
||||
clang::SourceManager* SM = Result.SourceManager;
|
||||
OS << sHipLaunchKernelGGL;
|
||||
if (caleeDecl->isTemplateInstantiation()) OS << "(";
|
||||
if (caleeDecl->isTemplateInstantiation()) OS << sHIP_KERNEL_NAME << "(";
|
||||
OS << readSourceText(*SM, calleeExpr->getSourceRange());
|
||||
if (caleeDecl->isTemplateInstantiation()) OS << ")";
|
||||
OS << ", ";
|
||||
|
||||
@@ -505,13 +505,19 @@ ihipDevice_t::ihipDevice_t(unsigned deviceId, unsigned deviceCnt, hc::accelerato
|
||||
: _deviceId(deviceId), _acc(acc), _state(0), _criticalData(this) {
|
||||
hsa_agent_t* agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
|
||||
if (agent) {
|
||||
int err = hsa_agent_get_info(
|
||||
int err;
|
||||
err = hsa_agent_get_info(
|
||||
*agent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_COMPUTE_UNIT_COUNT, &_computeUnits);
|
||||
if (err != HSA_STATUS_SUCCESS) {
|
||||
_computeUnits = 1;
|
||||
}
|
||||
|
||||
_hsaAgent = *agent;
|
||||
err = hsa_agent_get_info(
|
||||
*agent, (hsa_agent_info_t) HSA_AMD_AGENT_INFO_DRIVER_NODE_ID, &_driver_node_id);
|
||||
if (err != HSA_STATUS_SUCCESS){
|
||||
_driver_node_id = 0;
|
||||
}
|
||||
|
||||
_hsaAgent = *agent;
|
||||
} else {
|
||||
_hsaAgent.handle = static_cast<uint64_t>(-1);
|
||||
}
|
||||
|
||||
@@ -802,6 +802,9 @@ class ihipDevice_t {
|
||||
|
||||
// TODO - report this through device properties, base on HCC API call.
|
||||
int _isLargeBar;
|
||||
|
||||
// Node id reported by kfd for this device
|
||||
uint32_t _driver_node_id;
|
||||
|
||||
ihipCtx_t* _primaryCtx;
|
||||
|
||||
|
||||
@@ -19,7 +19,6 @@ 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 <hc_am.hpp>
|
||||
#include "hsa/hsa.h"
|
||||
#include "hsa/hsa_ext_amd.h"
|
||||
@@ -28,6 +27,8 @@ THE SOFTWARE.
|
||||
#include "hip_hcc_internal.h"
|
||||
#include "trace_helper.h"
|
||||
|
||||
#include <fstream>
|
||||
|
||||
__device__ char __hip_device_heap[__HIP_SIZE_OF_HEAP];
|
||||
__device__ uint32_t __hip_device_page_flag[__HIP_NUM_PAGES];
|
||||
|
||||
@@ -1967,16 +1968,27 @@ hipError_t hipMemGetInfo(size_t* free, size_t* total) {
|
||||
} else {
|
||||
e = hipErrorInvalidValue;
|
||||
}
|
||||
|
||||
|
||||
if (free) {
|
||||
// TODO - replace with kernel-level for reporting free memory:
|
||||
size_t deviceMemSize, hostMemSize, userMemSize;
|
||||
hc::am_memtracker_sizeinfo(device->_acc, &deviceMemSize, &hostMemSize, &userMemSize);
|
||||
|
||||
*free = device->_props.totalGlobalMem - deviceMemSize;
|
||||
|
||||
// Deduct the amount of memory from the free memory reported from the system
|
||||
if (HIP_HIDDEN_FREE_MEM) *free -= (size_t)HIP_HIDDEN_FREE_MEM * 1024 * 1024;
|
||||
if (!device->_driver_node_id) return ihipLogStatus(hipErrorInvalidDevice);
|
||||
|
||||
std::string fileName = std::string("/sys/class/kfd/kfd/topology/nodes/") + std::to_string(device->_driver_node_id) + std::string("/mem_banks/0/used_memory");
|
||||
std::ifstream file;
|
||||
file.open(fileName);
|
||||
if (!file) return ihipLogStatus(hipErrorFileNotFound);
|
||||
|
||||
std::string deviceSize;
|
||||
size_t deviceMemSize;
|
||||
|
||||
file >> deviceSize;
|
||||
file.close();
|
||||
if ((deviceMemSize=strtol(deviceSize.c_str(),NULL,10))){
|
||||
*free = device->_props.totalGlobalMem - deviceMemSize;
|
||||
// Deduct the amount of memory from the free memory reported from the system
|
||||
if (HIP_HIDDEN_FREE_MEM) *free -= (size_t)HIP_HIDDEN_FREE_MEM * 1024 * 1024;
|
||||
} else {
|
||||
return ihipLogStatus(hipErrorInvalidValue);
|
||||
}
|
||||
} else {
|
||||
e = hipErrorInvalidValue;
|
||||
}
|
||||
|
||||
@@ -0,0 +1,284 @@
|
||||
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args "-Xclang" "-fcuda-allow-variadic-functions"
|
||||
|
||||
/*
|
||||
Copyright (c) 2015-present 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.
|
||||
*/
|
||||
|
||||
// CHECK: #include "hip/hip_runtime.h"
|
||||
#include "cuda_runtime.h"
|
||||
|
||||
#include <algorithm>
|
||||
#include <stdlib.h>
|
||||
#include <stdio.h>
|
||||
#include <string.h>
|
||||
#include <math.h>
|
||||
#include <type_traits>
|
||||
|
||||
#define EXIT_WAIVED 2
|
||||
|
||||
const char* sampleName = "hipSimpleAtomicsTest";
|
||||
|
||||
using namespace std;
|
||||
// Auto-Verification Code
|
||||
bool testResult = true;
|
||||
|
||||
bool computeGoldBitwise(...) {
|
||||
return true;
|
||||
}
|
||||
|
||||
template<typename T, typename enable_if<is_integral<T>{}>::type* = nullptr>
|
||||
bool computeGoldBitwise(T* gpuData, int len) {
|
||||
T val = 0xff;
|
||||
for (int i = 0; i < len; ++i) {
|
||||
// 9th element should be 1
|
||||
val &= (2 * i + 7);
|
||||
}
|
||||
if (val != gpuData[8]) {
|
||||
printf("atomicAnd failed\n");
|
||||
return false;
|
||||
}
|
||||
|
||||
val = 0;
|
||||
for (int i = 0; i < len; ++i) {
|
||||
// 10th element should be 0xff
|
||||
val |= (1 << i);
|
||||
}
|
||||
if (val != gpuData[9]) {
|
||||
printf("atomicOr failed\n");
|
||||
return false;
|
||||
}
|
||||
|
||||
val = 0xff;
|
||||
for (int i = 0; i < len; ++i) {
|
||||
// 11th element should be 0xff
|
||||
val ^= i;
|
||||
}
|
||||
if (val != gpuData[10]) {
|
||||
printf("atomicXor failed\n");
|
||||
return false;
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
bool computeGold(T* gpuData, int len) {
|
||||
T val = 0;
|
||||
for (int i = 0; i < len; ++i) {
|
||||
val += 10;
|
||||
}
|
||||
if (val != gpuData[0]) {
|
||||
printf("atomicAdd failed\n");
|
||||
return false;
|
||||
}
|
||||
|
||||
val = 0;
|
||||
for (int i = 0; i < len; ++i) {
|
||||
val -= 10;
|
||||
}
|
||||
if (val != gpuData[1]) {
|
||||
printf("atomicSub failed\n");
|
||||
return false;
|
||||
}
|
||||
|
||||
bool found = false;
|
||||
for (T i = 0; i < len; ++i) {
|
||||
// third element should be a member of [0, len)
|
||||
if (i == gpuData[2]) {
|
||||
found = true;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (!found) {
|
||||
printf("atomicExch failed\n");
|
||||
return false;
|
||||
}
|
||||
|
||||
val = -(1 << 8);
|
||||
for (T i = 0; i < len; ++i) {
|
||||
// fourth element should be len-1
|
||||
val = max(val, i);
|
||||
}
|
||||
if (val != gpuData[3]) {
|
||||
printf("atomicMax failed\n");
|
||||
return false;
|
||||
}
|
||||
|
||||
val = 1 << 8;
|
||||
for (T i = 0; i < len; ++i) {
|
||||
val = min(val, i);
|
||||
}
|
||||
if (val != gpuData[4]) {
|
||||
printf("atomicMin failed\n");
|
||||
return false;
|
||||
}
|
||||
|
||||
int limit = 17;
|
||||
val = 0;
|
||||
for (int i = 0; i < len; ++i) {
|
||||
val = (val >= limit) ? 0 : val + 1;
|
||||
}
|
||||
if (val != gpuData[5]) {
|
||||
printf("atomicInc failed\n");
|
||||
return false;
|
||||
}
|
||||
|
||||
limit = 137;
|
||||
val = 0;
|
||||
for (int i = 0; i < len; ++i) {
|
||||
val = ((val == 0) || (val > limit)) ? limit : val - 1;
|
||||
}
|
||||
if (val != gpuData[6]) {
|
||||
printf("atomicDec failed\n");
|
||||
return false;
|
||||
}
|
||||
|
||||
found = false;
|
||||
for (T i = 0; i < len; ++i) {
|
||||
// eighth element should be a member of [0, len)
|
||||
if (i == gpuData[7]) {
|
||||
found = true;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (!found) {
|
||||
printf("atomicCAS failed\n");
|
||||
return false;
|
||||
}
|
||||
|
||||
return computeGoldBitwise(gpuData, len);
|
||||
}
|
||||
|
||||
__device__
|
||||
void testKernelExch(...) {}
|
||||
|
||||
template<typename T, typename enable_if<!is_same<T, double>{}>::type* = nullptr>
|
||||
__device__
|
||||
void testKernelExch(T* g_odata) {
|
||||
// access thread id
|
||||
const T tid = blockDim.x * blockIdx.x + threadIdx.x;
|
||||
// Atomic exchange
|
||||
atomicExch(&g_odata[2], tid);
|
||||
}
|
||||
|
||||
__device__
|
||||
void testKernelSub(...) {}
|
||||
|
||||
template<
|
||||
typename T,
|
||||
typename enable_if<
|
||||
is_same<T, int>{} || is_same<T, unsigned int>{}>::type* = nullptr>
|
||||
__device__
|
||||
void testKernelSub(T* g_odata) {
|
||||
// Atomic subtraction (final should be 0)
|
||||
atomicSub(&g_odata[1], 10);
|
||||
}
|
||||
|
||||
__device__
|
||||
void testKernelIntegral(...) {}
|
||||
|
||||
template<typename T, typename enable_if<is_integral<T>{}>::type* = nullptr>
|
||||
__device__
|
||||
void testKernelIntegral(T* g_odata) {
|
||||
// access thread id
|
||||
const T tid = blockDim.x * blockIdx.x + threadIdx.x;
|
||||
// Atomic maximum
|
||||
atomicMax(&g_odata[3], tid);
|
||||
// Atomic minimum
|
||||
atomicMin(&g_odata[4], tid);
|
||||
// Atomic increment (modulo 17+1)
|
||||
atomicInc((unsigned int*)&g_odata[5], 17);
|
||||
// Atomic decrement
|
||||
atomicDec((unsigned int*)&g_odata[6], 137);
|
||||
// Atomic compare-and-swap
|
||||
atomicCAS(&g_odata[7], tid - 1, tid);
|
||||
// Atomic AND
|
||||
atomicAnd(&g_odata[8], 2 * tid + 7);
|
||||
// Atomic OR
|
||||
atomicOr(&g_odata[9], 1 << tid);
|
||||
// Atomic XOR
|
||||
atomicXor(&g_odata[10], tid);
|
||||
testKernelSub(g_odata);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
__global__ void testKernel(T* g_odata) {
|
||||
// Atomic addition
|
||||
atomicAdd(&g_odata[0], 10);
|
||||
testKernelIntegral(g_odata);
|
||||
testKernelExch(g_odata);
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
void runTest() {
|
||||
// CHECK: hipDeviceProp_t deviceProp;
|
||||
cudaDeviceProp deviceProp;
|
||||
deviceProp.major = 0;
|
||||
deviceProp.minor = 0;
|
||||
int dev = 0;
|
||||
// CHECK: hipGetDeviceProperties(&deviceProp, dev);
|
||||
cudaGetDeviceProperties(&deviceProp, dev);
|
||||
// Statistics about the GPU device
|
||||
printf(
|
||||
"> GPU device has %d Multi-Processors, "
|
||||
"SM %d.%d compute capabilities\n\n",
|
||||
deviceProp.multiProcessorCount, deviceProp.major, deviceProp.minor);
|
||||
unsigned int numThreads = 256;
|
||||
unsigned int numBlocks = 64;
|
||||
unsigned int numData = 11;
|
||||
unsigned int memSize = sizeof(T) * numData;
|
||||
|
||||
// Allocate mem for the result on host side
|
||||
T* hOData = (T*)malloc(memSize);
|
||||
// Initialize the memory
|
||||
for (unsigned int i = 0; i < numData; i++) hOData[i] = 0;
|
||||
// To make the AND and XOR tests generate something other than 0...
|
||||
hOData[8] = hOData[10] = 0xff;
|
||||
// Allocate device memory for result
|
||||
T* dOData;
|
||||
// CHECK: hipMalloc((void**)&dOData, memSize);
|
||||
cudaMalloc((void**)&dOData, memSize);
|
||||
// Copy host memory to device to initialize to zero
|
||||
// CHECK: hipMemcpy(dOData, hOData, memSize, hipMemcpyHostToDevice);
|
||||
cudaMemcpy(dOData, hOData, memSize, cudaMemcpyHostToDevice);
|
||||
// Execute the kernel
|
||||
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(testKernel), dim3(numBlocks), dim3(numThreads), 0, 0, dOData);
|
||||
testKernel<<<numBlocks, numThreads>>>(dOData);
|
||||
// Copy result from device to host
|
||||
// CHECK: hipMemcpy(hOData, dOData, memSize, hipMemcpyDeviceToHost);
|
||||
cudaMemcpy(hOData, dOData, memSize, cudaMemcpyDeviceToHost);
|
||||
// Compute reference solution
|
||||
testResult = computeGold(hOData, numThreads * numBlocks);
|
||||
// Cleanup memory
|
||||
free(hOData);
|
||||
// CHECK: hipFree(dOData);
|
||||
cudaFree(dOData);
|
||||
}
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
printf("%s starting...\n", sampleName);
|
||||
runTest<int>();
|
||||
runTest<unsigned int>();
|
||||
runTest<unsigned long long>();
|
||||
runTest<float>();
|
||||
runTest<double>();
|
||||
// CHECK: hipDeviceReset();
|
||||
cudaDeviceReset();
|
||||
printf("%s completed, returned %s\n", sampleName, testResult ? "OK" : "ERROR!");
|
||||
exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE);
|
||||
}
|
||||
+5
-5
@@ -121,7 +121,7 @@ struct runner
|
||||
const size_t states_size = blocks * threads;
|
||||
// CHECK: CUDA_CALL(hipMalloc((void **)&states, states_size * sizeof(GeneratorState)));
|
||||
CUDA_CALL(cudaMalloc((void **)&states, states_size * sizeof(GeneratorState)));
|
||||
// CHECK: hipLaunchKernelGGL((init_kernel), dim3(blocks), dim3(threads), 0, 0, states, seed, offset);
|
||||
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(init_kernel), dim3(blocks), dim3(threads), 0, 0, states, seed, offset);
|
||||
init_kernel<<<blocks, threads>>>(states, seed, offset);
|
||||
// CHECK: CUDA_CALL(hipPeekAtLastError());
|
||||
// CHECK: CUDA_CALL(hipDeviceSynchronize());
|
||||
@@ -142,7 +142,7 @@ struct runner
|
||||
const GenerateFunc& generate_func,
|
||||
const Extra extra)
|
||||
{
|
||||
// CHECK: hipLaunchKernelGGL((generate_kernel), dim3(blocks), dim3(threads), 0, 0, states, data, size, generate_func, extra);
|
||||
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(generate_kernel), dim3(blocks), dim3(threads), 0, 0, states, data, size, generate_func, extra);
|
||||
generate_kernel<<<blocks, threads>>>(states, data, size, generate_func, extra);
|
||||
}
|
||||
};
|
||||
@@ -223,7 +223,7 @@ struct runner<curandStateMtgp32_t>
|
||||
const GenerateFunc& generate_func,
|
||||
const Extra extra)
|
||||
{
|
||||
// CHECK: hipLaunchKernelGGL((generate_kernel), dim3(std::min((size_t)200, blocks)), dim3(256), 0, 0, states, data, size, generate_func, extra);
|
||||
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(generate_kernel), dim3(std::min((size_t)200, blocks)), dim3(256), 0, 0, states, data, size, generate_func, extra);
|
||||
generate_kernel<<<std::min((size_t)200, blocks), 256>>>(states, data, size, generate_func, extra);
|
||||
}
|
||||
};
|
||||
@@ -304,7 +304,7 @@ struct runner<curandStateSobol32_t>
|
||||
CUDA_CALL(cudaMemcpy(directions, h_directions, size, cudaMemcpyHostToDevice));
|
||||
|
||||
const size_t blocks_x = next_power2((blocks + dimensions - 1) / dimensions);
|
||||
// CHECK: hipLaunchKernelGGL((init_kernel), dim3(blocks_x, dimensions), dim3(threads), 0, 0, states, directions, offset);
|
||||
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(init_kernel), dim3(blocks_x, dimensions), dim3(threads), 0, 0, states, directions, offset);
|
||||
init_kernel<<<dim3(blocks_x, dimensions), threads>>>(states, directions, offset);
|
||||
// CHECK: CUDA_CALL(hipPeekAtLastError());
|
||||
// CHECK: CUDA_CALL(hipDeviceSynchronize());
|
||||
@@ -329,7 +329,7 @@ struct runner<curandStateSobol32_t>
|
||||
const Extra extra)
|
||||
{
|
||||
const size_t blocks_x = next_power2((blocks + dimensions - 1) / dimensions);
|
||||
// CHECK: hipLaunchKernelGGL((generate_kernel), dim3(blocks_x, dimensions), dim3(threads), 0, 0, states, data, size / dimensions, generate_func, extra);
|
||||
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(generate_kernel), dim3(blocks_x, dimensions), dim3(threads), 0, 0, states, data, size / dimensions, generate_func, extra);
|
||||
generate_kernel<<<dim3(blocks_x, dimensions), threads>>>(states, data, size / dimensions, generate_func, extra);
|
||||
}
|
||||
};
|
||||
|
||||
@@ -10,7 +10,7 @@
|
||||
#define KERNEL_CALL_AS_MACRO axpy<float><<<1, kDataLen>>>
|
||||
#define KERNEL_NAME_MACRO axpy<float>
|
||||
|
||||
// CHECK: #define COMPLETE_LAUNCH hipLaunchKernelGGL((axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y)
|
||||
// CHECK: #define COMPLETE_LAUNCH hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y)
|
||||
#define COMPLETE_LAUNCH axpy<<<1, kDataLen>>>(a, device_x, device_y)
|
||||
|
||||
|
||||
@@ -48,22 +48,22 @@ int main(int argc, char* argv[]) {
|
||||
cudaMemcpy(device_x, host_x, kDataLen * sizeof(float), cudaMemcpyHostToDevice);
|
||||
|
||||
// Launch the kernel in numerous different strange ways to exercise the prerocessor.
|
||||
// CHECK: hipLaunchKernelGGL((axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
|
||||
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
|
||||
axpy<<<1, kDataLen>>>(a, device_x, device_y);
|
||||
|
||||
// CHECK: hipLaunchKernelGGL((axpy<float>), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
|
||||
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy<float>), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
|
||||
axpy<float><<<1, kDataLen>>>(a, device_x, device_y);
|
||||
|
||||
// CHECK: hipLaunchKernelGGL((axpy<float>), dim3(1), dim3(kDataLen), 0, 0, a, TOKEN_PASTE(device, _x), device_y);
|
||||
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy<float>), dim3(1), dim3(kDataLen), 0, 0, a, TOKEN_PASTE(device, _x), device_y);
|
||||
axpy<float><<<1, kDataLen>>>(a, TOKEN_PASTE(device, _x), device_y);
|
||||
|
||||
// CHECK: hipLaunchKernelGGL((axpy<float>), dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
|
||||
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy<float>), dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
|
||||
axpy<float><<<1, kDataLen>>>(ARG_LIST_AS_MACRO);
|
||||
|
||||
// CHECK: hipLaunchKernelGGL((KERNEL_NAME_MACRO), dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
|
||||
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(KERNEL_NAME_MACRO), dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
|
||||
KERNEL_NAME_MACRO<<<1, kDataLen>>>(ARG_LIST_AS_MACRO);
|
||||
|
||||
// CHECK: hipLaunchKernelGGL((axpy<float>), dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
|
||||
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy<float>), dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
|
||||
KERNEL_CALL_AS_MACRO(ARG_LIST_AS_MACRO);
|
||||
|
||||
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), 0, 0);
|
||||
|
||||
@@ -94,7 +94,7 @@ int main(int argc, char *argv[])
|
||||
const unsigned threadsPerBlock = 256;
|
||||
|
||||
printf ("info: launch 'vector_square' kernel\n");
|
||||
// CHECK: hipLaunchKernelGGL((vector_square), dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
|
||||
// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(vector_square), dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N);
|
||||
vector_square <<<blocks, threadsPerBlock>>> (C_d, A_d, N);
|
||||
|
||||
printf ("info: copy Device2Host\n");
|
||||
|
||||
@@ -27,21 +27,10 @@
|
||||
*/
|
||||
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "../test_common.h"
|
||||
#include <stdio.h>
|
||||
#include "LaunchKernel.h"
|
||||
|
||||
#define HIPCHECK(error) \
|
||||
{ \
|
||||
hipError_t localError = error; \
|
||||
if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \
|
||||
printf("%serror: '%s'(%d) from %s at %s:%d%s\n", "\x1B[31m", hipGetErrorString(localError), \
|
||||
localError, #error, __FILE__, __LINE__, "\x1B[0m"); \
|
||||
return false; \
|
||||
} \
|
||||
}
|
||||
|
||||
|
||||
bool LaunchKernelArg()
|
||||
{
|
||||
dim3 blocks = {1,1,1};
|
||||
|
||||
@@ -17,13 +17,23 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <iostream>
|
||||
#include <iomanip>
|
||||
#if __CUDACC__
|
||||
#include <sys/time.h>
|
||||
#else
|
||||
#include <chrono>
|
||||
/*
|
||||
* File is intended to C and CPP compliant hence any CPP specic changes
|
||||
* should be added into CPP section
|
||||
*
|
||||
*/
|
||||
|
||||
#ifdef __cplusplus
|
||||
#include <iostream>
|
||||
#include <iomanip>
|
||||
#if __CUDACC__
|
||||
#include <sys/time.h>
|
||||
#else
|
||||
#include <chrono>
|
||||
#endif
|
||||
#endif
|
||||
|
||||
// ************************ GCC section **************************
|
||||
#include <stddef.h>
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
@@ -41,14 +51,6 @@ THE SOFTWARE.
|
||||
#define KCYN "\x1B[36m"
|
||||
#define KWHT "\x1B[37m"
|
||||
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC
|
||||
#define TYPENAME(T) typeid(T).name()
|
||||
#else
|
||||
#define TYPENAME(T) "?"
|
||||
#endif
|
||||
|
||||
|
||||
#define passed() \
|
||||
printf("%sPASSED!%s\n", KGRN, KNRM); \
|
||||
exit(0);
|
||||
@@ -114,6 +116,15 @@ extern int p_gpuDevice;
|
||||
extern unsigned p_verbose;
|
||||
extern int p_tests;
|
||||
|
||||
// ********************* CPP section *********************
|
||||
#ifdef __cplusplus
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC
|
||||
#define TYPENAME(T) typeid(T).name()
|
||||
#else
|
||||
#define TYPENAME(T) "?"
|
||||
#endif
|
||||
|
||||
namespace HipTest {
|
||||
|
||||
// Returns the current system time in microseconds
|
||||
@@ -470,3 +481,4 @@ struct MemTraits<MemcpyAsync> {
|
||||
};
|
||||
|
||||
}; // namespace HipTest
|
||||
#endif //__cplusplus
|
||||
|
||||
Reference in New Issue
Block a user