AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

Merge pull request #1486 from emankov/hipify-clang

Telusuri Sumber 
[HIPIFY][perl] Generate transformExternShared function
This commit is contained in:
Evgeny Mankov
2019-09-30 17:51:07 +03:00
committed by GitHub
orang tua 0e91fb36d2 bf5546f35e
melakukan fae82d73d7
2 mengubah file dengan 25 tambahan dan 26 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/bin/hipify-perl
+13 -24
Melihat File
@@ -91,7 +91,7 @@ sub printStats {
printf STDERR "%s:%d ", $stat, $counts{$stat};
}
printf STDERR ")\n warn:%d LOC:%d", $warnings, $loc;
};
}
sub addStats {
my $dest_ref = shift();
@@ -99,7 +99,7 @@ sub addStats {
foreach $key (keys %adder) {
$dest_ref->{$key} += $adder{$key};
}
};
}
sub clearStats {
my $dest_ref = shift();
@@ -107,7 +107,7 @@ sub clearStats {
foreach $stat(@statNames) {
$dest_ref->{$stat} = 0;
}
};
}
sub simpleSubstitutions {
$ft{'error'} += s/\bcudaGetErrorName\b/hipGetErrorName/g;
@@ -1594,6 +1594,14 @@ sub simpleSubstitutions {
$ft{'define'} += s/\bcudaTextureTypeCubemapLayered\b/hipTextureTypeCubemapLayered/g;
}
# CUDA extern __shared__ syntax replace with HIP_DYNAMIC_SHARED() macro
sub transformExternShared {
no warnings qw/uninitialized/;
my $k = 0;
$k += s/extern\s+([\w\(\)]+)?\s*__shared__\s+([\w:<>\s]+)\s+(\w+)\s*\[\s*\]\s*;/HIP_DYNAMIC_SHARED($1 $2, $3)/g;
$ft{ 'extern_shared' } += $k;
}
# Count of transforms in all files
my %tt;
clearStats(\%tt, \@statNames);
@@ -1643,27 +1651,8 @@ while (@ARGV) {
$countKeywords += m/__global__/;
$countKeywords += m/__shared__/;
# CUDA extern __shared__ syntax
# Note these only work if declaration is on a single line.
{
# Match uses ? for <.*> which will be unitialized if this is not present in launch syntax
no warnings qw/uninitialized/;
my $k = 0;
# Replace as HIP_DYNAMIC_SHARED() macro
# Match patterns for the below regular expression:
#'extern __shared__ double foo[];'
#'extern __shared__ unsigned int foo[];'
#'extern volatile __shared__ double foo[];'
#'extern volatile __shared__ unsigned int sdata[];'
#'extern __shared__ volatile unsigned int sdata[];'
#'extern __shared__ T s[];'
#'extern __shared__ T::type s[];'
#'extern __shared__ blah<T>::type s[];'
#'extern __shared__ typename mapper<Float>::type s_data[];'
#'extern __attribute__((used)) __shared__ typename mapper<Float>::type s_data[];'
$k += s/extern\s+([\w\(\)]+)?\s*__shared__\s+([\w:<>\s]+)\s+(\w+)\s*\[\s*\]\s*;/HIP_DYNAMIC_SHARED($1 $2, $3)/g;
$ft{'extern_shared'} += $k;
}
transformExternShared();
# CUDA Launch Syntax. Note these only work if launch is on a single line.
{
# Match uses ? for <.*> which will be unitialized if this is not present in launch syntax
hipamd/hipify-clang/src/CUDA2HIP_Perl.cpp
+12 -2
Melihat File
@@ -143,13 +143,13 @@ 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;
}
}
@@ -179,6 +179,15 @@ namespace perl {
*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 << "no warnings qw/uninitialized/;" << std::endl;
*streamPtr.get() << tab << "my $k = 0;" << 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 generateHostFunctions(std::unique_ptr<std::ostream>& streamPtr) {
*streamPtr.get() << std::endl << sSub << " transformHostFunctions" << " {" << std::endl << tab << sMy;
std::set<std::string> &funcSet = DeviceSymbolFunctions0;
@@ -295,6 +304,7 @@ namespace perl {
*streamPtr.get() << "\"" << counterNames[NUM_CONV_TYPES - 1] << "\");" << std::endl;
generateStatFunctions(streamPtr);
generateSimpleSubstitutions(streamPtr);
generateExternShared(streamPtr);
generateHostFunctions(streamPtr);
generateDeviceFunctions(streamPtr);
streamPtr.get()->flush();