diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl index 6ec8f7941a..4f926a779c 100755 --- a/projects/hip/bin/hipify-perl +++ b/projects/hip/bin/hipify-perl @@ -1757,7 +1757,7 @@ while (@ARGV) { $ft{'kernel_func'} += countSupportedDeviceFunctions(); } - $ft{'memory'} += transformSymbolFunctions(); + transformHostFunctions(); # Print it! # TODO - would like to move this code outside loop but it uses $_ which contains the whole file. @@ -1813,7 +1813,7 @@ if ($count_conversions) { } } -sub transformSymbolFunctions +sub transformHostFunctions { my $m = 0; foreach $func ( @@ -1832,6 +1832,18 @@ sub transformSymbolFunctions { $m += s/(?\($2\),/g + } + foreach $func ( + "hipFuncGetAttributes" + ) + { + $m += s/(?\($3\)$4/g; + } return $m; } diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp index 26a72886c7..5407cae35c 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Perl.cpp @@ -45,12 +45,19 @@ namespace perl { const std::string sForeach = "foreach $func (\n"; const std::string sMy = "my $m = 0;\n"; - void generateSymbolFunctions(std::unique_ptr& perlStreamPtr) { - *perlStreamPtr.get() << "\n" << sSub << " transformSymbolFunctions\n" << "{\n" << tab << sMy; + void generateHostFunctions(std::unique_ptr& perlStreamPtr) { + *perlStreamPtr.get() << "\n" << sSub << " transformHostFunctions\n" << "{\n" << tab << sMy; std::set &funcSet = DeviceSymbolFunctions0; - for (int i = 0; i < 2; ++i) { + const std::string s0 = "$m += s/(?flush(); bool ret = true; diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Scripting.h b/projects/hip/hipify-clang/src/CUDA2HIP_Scripting.h index 9469f25845..76b103735b 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Scripting.h +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Scripting.h @@ -24,6 +24,11 @@ THE SOFTWARE. extern std::set DeviceSymbolFunctions0; extern std::set DeviceSymbolFunctions1; +extern std::set ReinterpretFunctions0; +extern std::set ReinterpretFunctions1; + +extern std::string sHIP_SYMBOL; +extern std::string s_reinterpret_cast; namespace perl { diff --git a/projects/hip/hipify-clang/src/HipifyAction.cpp b/projects/hip/hipify-clang/src/HipifyAction.cpp index a2380db139..e331cfbec5 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.cpp +++ b/projects/hip/hipify-clang/src/HipifyAction.cpp @@ -37,8 +37,8 @@ namespace ct = clang::tooling; namespace mat = clang::ast_matchers; const std::string sHIP_DYNAMIC_SHARED = "HIP_DYNAMIC_SHARED"; -const std::string sHIP_SYMBOL = "HIP_SYMBOL"; -const std::string s_reinterpret_cast = "reinterpret_cast"; +std::string sHIP_SYMBOL = "HIP_SYMBOL"; +std::string s_reinterpret_cast = "reinterpret_cast"; const std::string sHipLaunchKernelGGL = "hipLaunchKernelGGL("; const std::string sDim3 = "dim3("; @@ -68,6 +68,14 @@ std::set ReinterpretFunctions{ {sCudaFuncGetAttributes} }; +std::set ReinterpretFunctions0{ + {sCudaFuncSetCacheConfig} +}; + +std::set ReinterpretFunctions1{ + {sCudaFuncGetAttributes} +}; + void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) { clang::SourceManager& SM = getCompilerInstance().getSourceManager(); size_t begin = 0; diff --git a/projects/hip/tests/hipify-clang/unit_tests/samples/reinterpret_cast.cu b/projects/hip/tests/hipify-clang/unit_tests/casts/reinterpret_cast.cu similarity index 96% rename from projects/hip/tests/hipify-clang/unit_tests/samples/reinterpret_cast.cu rename to projects/hip/tests/hipify-clang/unit_tests/casts/reinterpret_cast.cu index fe67629ec1..6b6f4f5dde 100644 --- a/projects/hip/tests/hipify-clang/unit_tests/samples/reinterpret_cast.cu +++ b/projects/hip/tests/hipify-clang/unit_tests/casts/reinterpret_cast.cu @@ -22,7 +22,6 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include // CHECK: #include #include @@ -32,7 +31,7 @@ void fn(float* px, float* py) { __shared__ double b[69]; for (auto&& x : b) x = *py++; for (auto&& x : a) x = *px++ > 0.0; - for (auto&& x : a) if (x)* --py = *--px; + for (auto&& x : a) if (x) *--py = *--px; } int main() {