diff --git a/bin/hipify-perl b/bin/hipify-perl index 58001988..46349130 100755 --- a/bin/hipify-perl +++ b/bin/hipify-perl @@ -5619,7 +5619,7 @@ sub simpleSubstitutions { subst("curand_mtgp32.h", "hiprand\/hiprand_kernel.h", "include"); subst("curand_mtgp32_host.h", "hiprand\/hiprand_mtgp32_host.h", "include"); subst("curand_mtgp32_kernel.h", "hiprand\/hiprand_kernel.h", "include"); - subst("curand_mtgp32dc_p_11213.h", "rocrand_mtgp32_11213.h", "include"); + subst("curand_mtgp32dc_p_11213.h", "rocrand\/rocrand_mtgp32_11213.h", "include"); subst("curand_normal.h", "hiprand\/hiprand_kernel.h", "include"); subst("curand_normal_static.h", "hiprand\/hiprand_kernel.h", "include"); subst("curand_philox4x32_x.h", "hiprand\/hiprand_kernel.h", "include"); diff --git a/src/CUDA2HIP.cpp b/src/CUDA2HIP.cpp index 11d7395b..d6d8feb0 100644 --- a/src/CUDA2HIP.cpp +++ b/src/CUDA2HIP.cpp @@ -48,23 +48,23 @@ const std::map CUDA_INCLUDE_MAP { {"cublas_api.h", {"hipblas.h", "rocblas.h", CONV_INCLUDE, API_BLAS, 0}}, {"cublasLt.h", {"hipblaslt.h", "", CONV_INCLUDE, API_BLAS, 0}}, // cuRAND includes - {"curand.h", {"hiprand/hiprand.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_RAND, 0}}, - {"curand_kernel.h", {"hiprand/hiprand_kernel.h", "", CONV_INCLUDE, API_RAND, 0}}, - {"curand_discrete.h", {"hiprand/hiprand_kernel.h", "", CONV_INCLUDE, API_RAND, 0}}, - {"curand_discrete2.h", {"hiprand/hiprand_kernel.h", "", CONV_INCLUDE, API_RAND, 0}}, - {"curand_globals.h", {"hiprand/hiprand_kernel.h", "", CONV_INCLUDE, API_RAND, 0}}, - {"curand_lognormal.h", {"hiprand/hiprand_kernel.h", "", CONV_INCLUDE, API_RAND, 0}}, - {"curand_mrg32k3a.h", {"hiprand/hiprand_kernel.h", "", CONV_INCLUDE, API_RAND, 0}}, - {"curand_mtgp32.h", {"hiprand/hiprand_kernel.h", "", CONV_INCLUDE, API_RAND, 0}}, - {"curand_mtgp32_host.h", {"hiprand/hiprand_mtgp32_host.h", "", CONV_INCLUDE, API_RAND, 0}}, - {"curand_mtgp32_kernel.h", {"hiprand/hiprand_kernel.h", "", CONV_INCLUDE, API_RAND, 0}}, - {"curand_mtgp32dc_p_11213.h", {"rocrand_mtgp32_11213.h", "", CONV_INCLUDE, API_RAND, 0}}, - {"curand_normal.h", {"hiprand/hiprand_kernel.h", "", CONV_INCLUDE, API_RAND, 0}}, - {"curand_normal_static.h", {"hiprand/hiprand_kernel.h", "", CONV_INCLUDE, API_RAND, 0}}, - {"curand_philox4x32_x.h", {"hiprand/hiprand_kernel.h", "", CONV_INCLUDE, API_RAND, 0}}, - {"curand_poisson.h", {"hiprand/hiprand_kernel.h", "", CONV_INCLUDE, API_RAND, 0}}, - {"curand_precalc.h", {"hiprand/hiprand_kernel.h", "", CONV_INCLUDE, API_RAND, 0}}, - {"curand_uniform.h", {"hiprand/hiprand_kernel.h", "", CONV_INCLUDE, API_RAND, 0}}, + {"curand.h", {"hiprand/hiprand.h", "rocrand/rocrand.h", CONV_INCLUDE_CUDA_MAIN_H, API_RAND, 0}}, + {"curand_kernel.h", {"hiprand/hiprand_kernel.h", "rocrand/rocrand_kernel.h", CONV_INCLUDE, API_RAND, 0}}, + {"curand_discrete.h", {"hiprand/hiprand_kernel.h", "rocrand/rocrand_discrete.h", CONV_INCLUDE, API_RAND, 0}}, + {"curand_discrete2.h", {"hiprand/hiprand_kernel.h", "rocrand/rocrand_discrete.h", CONV_INCLUDE, API_RAND, 0}}, + {"curand_globals.h", {"hiprand/hiprand_kernel.h", "rocrand/rocrand_common.h", CONV_INCLUDE, API_RAND, 0}}, + {"curand_lognormal.h", {"hiprand/hiprand_kernel.h", "rocrand/rocrand_log_normal.h", CONV_INCLUDE, API_RAND, 0}}, + {"curand_mrg32k3a.h", {"hiprand/hiprand_kernel.h", "rocrand/rocrand_mrg32k3a.h", CONV_INCLUDE, API_RAND, 0}}, + {"curand_mtgp32.h", {"hiprand/hiprand_kernel.h", "rocrand/rocrand_mtgp32.h", CONV_INCLUDE, API_RAND, 0}}, + {"curand_mtgp32_host.h", {"hiprand/hiprand_mtgp32_host.h", "rocrand/rocrand_mtgp32.h", CONV_INCLUDE, API_RAND, 0}}, + {"curand_mtgp32_kernel.h", {"hiprand/hiprand_kernel.h", "rocrand/rocrand_mtgp32.h", CONV_INCLUDE, API_RAND, 0}}, + {"curand_mtgp32dc_p_11213.h", {"rocrand/rocrand_mtgp32_11213.h", "rocrand/rocrand_mtgp32_11213.h", CONV_INCLUDE, API_RAND, 0}}, + {"curand_normal.h", {"hiprand/hiprand_kernel.h", "rocrand/rocrand_normal.h", CONV_INCLUDE, API_RAND, 0}}, + {"curand_normal_static.h", {"hiprand/hiprand_kernel.h", "rocrand/rocrand_normal.h", CONV_INCLUDE, API_RAND, 0}}, + {"curand_philox4x32_x.h", {"hiprand/hiprand_kernel.h", "rocrand/rocrand_philox4x32_10.h", CONV_INCLUDE, API_RAND, 0}}, + {"curand_poisson.h", {"hiprand/hiprand_kernel.h", "rocrand/rocrand_poisson.h", CONV_INCLUDE, API_RAND, 0}}, + {"curand_precalc.h", {"hiprand/hiprand_kernel.h", "rocrand/rocrand_xorwow_precomputed.h", CONV_INCLUDE, API_RAND, 0}}, + {"curand_uniform.h", {"hiprand/hiprand_kernel.h", "rocrand/rocrand_uniform.h", CONV_INCLUDE, API_RAND, 0}}, // cuDNN includes {"cudnn.h", {"hipDNN.h", "miopen/miopen.h", CONV_INCLUDE_CUDA_MAIN_H, API_DNN, 0}}, // cuFFT includes diff --git a/src/HipifyAction.cpp b/src/HipifyAction.cpp index 9258bca9..65149193 100644 --- a/src/HipifyAction.cpp +++ b/src/HipifyAction.cpp @@ -51,6 +51,9 @@ std::string s_int64_t = "int64_t"; const std::string sHipLaunchKernelGGL = "hipLaunchKernelGGL"; const std::string sDim3 = "dim3("; const std::string s_hiprand_kernel_h = "hiprand/hiprand_kernel.h"; +const std::string s_rocrand_discrete_t = "rocrand/rocrand_discrete.h"; +const std::string s_rocrand_mtgp32_t = "rocrand/rocrand_mtgp32.h"; +const std::string s_rocrand_normal_t = "rocrand/rocrand_normal.h"; const std::string s_hiprand_h = "hiprand/hiprand.h"; const std::string sOnce = "once"; const std::string s_string_literal = "[string literal]"; @@ -2383,15 +2386,18 @@ bool HipifyAction::Exclude(const hipCounter &hipToken) { insertedBLASHeader = true; return false; case API_RAND: - if (hipToken.hipName == s_hiprand_kernel_h) { - if (insertedRAND_kernelHeader) return true; - insertedRAND_kernelHeader = true; - return false; - } else if (hipToken.hipName == s_hiprand_h) { - if (insertedRANDHeader) return true; - insertedRANDHeader = true; - return false; + if (!Statistics::isToRoc(hipToken)) { + if (hipToken.hipName == s_hiprand_kernel_h) { + if (insertedRAND_kernelHeader) return true; + insertedRAND_kernelHeader = true; + return false; + } else if (hipToken.hipName == s_hiprand_h) { + if (insertedRANDHeader) return true; + insertedRANDHeader = true; + return false; + } } + return false; case API_DNN: if (insertedDNNHeader) return true; insertedDNNHeader = true; @@ -2436,9 +2442,23 @@ bool HipifyAction::Exclude(const hipCounter &hipToken) { if (hipToken.hipName.empty()) return true; switch (hipToken.apiType) { case API_RAND: - if (hipToken.hipName == s_hiprand_kernel_h) { - if (insertedRAND_kernelHeader) return true; - insertedRAND_kernelHeader = true; + if (Statistics::isToRoc(hipToken)) { + if (hipToken.rocName == s_rocrand_discrete_t) { + if (insertedRocRAND_discreteHeader) return true; + insertedRocRAND_discreteHeader = true; + } else if (hipToken.rocName == s_rocrand_mtgp32_t) { + if (insertedRocRAND_mtgp32Header) return true; + insertedRocRAND_mtgp32Header = true; + } else if (hipToken.rocName == s_rocrand_normal_t) { + if (insertedRocRAND_normalHeader) return true; + insertedRocRAND_normalHeader = true; + } + } + else { + if (hipToken.hipName == s_hiprand_kernel_h) { + if (insertedRAND_kernelHeader) return true; + insertedRAND_kernelHeader = true; + } } return false; default: diff --git a/src/HipifyAction.h b/src/HipifyAction.h index 55bc77bb..a88dd813 100644 --- a/src/HipifyAction.h +++ b/src/HipifyAction.h @@ -53,6 +53,9 @@ class HipifyAction : public clang::ASTFrontendAction, bool insertedBLASHeader_V2 = false; bool insertedRANDHeader = false; bool insertedRAND_kernelHeader = false; + bool insertedRocRAND_discreteHeader = false; + bool insertedRocRAND_mtgp32Header = false; + bool insertedRocRAND_normalHeader = false; bool insertedDNNHeader = false; bool insertedFFTHeader = false; bool insertedSPARSEHeader = false; diff --git a/tests/unit_tests/headers/headers_test_09.cu b/tests/unit_tests/headers/headers_test_09.cu index 078645e2..9a01dee5 100644 --- a/tests/unit_tests/headers/headers_test_09.cu +++ b/tests/unit_tests/headers/headers_test_09.cu @@ -47,7 +47,7 @@ // CHECK-NOT: #include "curand_uniform.h" // CHECK: #include "hiprand/hiprand_mtgp32_host.h" -// CHECK: #include "rocrand_mtgp32_11213.h" +// CHECK: #include "rocrand/rocrand_mtgp32_11213.h" // CHECK: #include diff --git a/tests/unit_tests/headers/headers_test_09_12000.cu b/tests/unit_tests/headers/headers_test_09_12000.cu index 823b8346..f0bc47d8 100644 --- a/tests/unit_tests/headers/headers_test_09_12000.cu +++ b/tests/unit_tests/headers/headers_test_09_12000.cu @@ -46,7 +46,7 @@ // CHECK-NOT: #include "curand_uniform.h" // CHECK: #include "hiprand/hiprand_mtgp32_host.h" -// CHECK: #include "rocrand_mtgp32_11213.h" +// CHECK: #include "rocrand/rocrand_mtgp32_11213.h" // CHECK: #include diff --git a/tests/unit_tests/headers/headers_test_09_rocrand.cu b/tests/unit_tests/headers/headers_test_09_rocrand.cu new file mode 100644 index 00000000..ca674d29 --- /dev/null +++ b/tests/unit_tests/headers/headers_test_09_rocrand.cu @@ -0,0 +1,112 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --roc %clang_args + +// CHECK: #include +// CHECK: #include + +// CHECK-NOT: #include +// CHECK-NOT: #include + +// CHECK: #include "hip/hip_runtime_api.h" +// CHECK: #include "hip/channel_descriptor.h" +// CHECK: #include "hip/device_functions.h" +// CHECK: #include "hip/driver_types.h" +// CHECK: #include "hip/hip_complex.h" +// CHECK: #include "hip/hip_fp16.h" +// CHECK: #include "hip/hip_texture_types.h" +// CHECK: #include "hip/hip_vector_types.h" + +// CHECK: #include + +// CHECK: #include + +// CHECK: #include "rocrand/rocrand.h" +// CHECK: #include "rocrand/rocrand_kernel.h" + +// CHECK: #include + +// CHECK: #include "rocrand/rocrand_discrete.h" +// CHECK: #include "rocrand/rocrand_common.h" +// CHECK: #include "rocrand/rocrand_log_normal.h" +// CHECK: #include "rocrand/rocrand_mrg32k3a.h" +// CHECK: #include "rocrand/rocrand_mtgp32.h" +// CHECK: #include "rocrand/rocrand_mtgp32_11213.h" +// CHECK: #include "rocrand/rocrand_normal.h" +// CHECK: #include "rocrand/rocrand_philox4x32_10.h" +// CHECK: #include "rocrand/rocrand_poisson.h" +// CHECK: #include "rocrand/rocrand_xorwow_precomputed.h" +// CHECK: #include "rocrand/rocrand_uniform.h" + +// CHECK-NOT: #include "rocrand/rocrand.h" +// CHECK-NOT: #include "rocrand/rocrand_kernel.h" +// CHECK-NOT: #include "rocrand/rocrand_discrete.h" +// CHECK-NOT: #include "rocrand/rocrand_mtgp32.h" +// CHECK-NOT: #include "rocrand/rocrand_normal.h" + +// CHECK-NOT: #include "curand_discrete.h" +// CHECK-NOT: #include "curand_discrete2.h" +// CHECK-NOT: #include "curand_globals.h" +// CHECK-NOT: #include "curand_lognormal.h" +// CHECK-NOT: #include "curand_mrg32k3a.h" +// CHECK-NOT: #include "curand_mtgp32.h" +// CHECK-NOT: #include "curand_mtgp32_host.h" +// CHECK-NOT: #include "curand_mtgp32_kernel.h" +// CHECK-NOT: #include "curand_mtgp32dc_p_11213.h" +// CHECK-NOT: #include "curand_normal.h" +// CHECK-NOT: #include "curand_normal_static.h" +// CHECK-NOT: #include "curand_philox4x32_x.h" +// CHECK-NOT: #include "curand_poisson.h" +// CHECK-NOT: #include "curand_precalc.h" +// CHECK-NOT: #include "curand_uniform.h" + +// CHECK: #include + +// CHECK: #include "hipfft/hipfft.h" +// CHECK: #include "rocsparse.h" + +#include +// CHECK-NOT: #include + +#include + +#include +// CHECK-NOT: #include + +#include "cuda_runtime_api.h" +#include "channel_descriptor.h" +#include "device_functions.h" +#include "driver_types.h" +#include "cuComplex.h" +#include "cuda_fp16.h" +#include "cuda_texture_types.h" +#include "vector_types.h" + +#include + +#include + +#include "curand.h" +#include "curand_kernel.h" + +#include + +#include "curand_discrete.h" +#include "curand_discrete2.h" +#include "curand_globals.h" +#include "curand_lognormal.h" +#include "curand_mrg32k3a.h" +#include "curand_mtgp32.h" +#include "curand_mtgp32_host.h" +#include "curand_mtgp32_kernel.h" +#include "curand_mtgp32dc_p_11213.h" +#include "curand_normal.h" +#include "curand_normal_static.h" +#include "curand_philox4x32_x.h" +#include "curand_poisson.h" +#include "curand_precalc.h" +#include "curand_uniform.h" + +#include + +#include "cufft.h" + +#include "cusparse.h" diff --git a/tests/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp b/tests/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp index 6eb492d0..f69b2d1e 100644 --- a/tests/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp +++ b/tests/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp @@ -39,7 +39,7 @@ #include // CHECK: #include #include -// CHECK: #include +// CHECK: #include #include // CHECK: if ((x) != hipSuccess) {