diff --git a/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md b/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md index c35eeb26b6..0561dd6144 100644 --- a/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md +++ b/hipamd/docs/markdown/CURAND_API_supported_by_HIP.md @@ -134,6 +134,8 @@ | `curandSetPseudoRandomGeneratorSeed` | `hiprandSetPseudoRandomGeneratorSeed` | | `curandSetQuasiRandomGeneratorDimensions` | `hiprandSetQuasiRandomGeneratorDimensions` | | `curandSetStream` | `hiprandSetStream` | +| `curandMakeMTGP32Constants` | `hiprandMakeMTGP32Constants` | +| `curandMakeMTGP32KernelState` | `hiprandMakeMTGP32KernelState` | ## **3. Device API Functions** diff --git a/hipamd/hipify-clang/src/CUDA2HIP.cpp b/hipamd/hipify-clang/src/CUDA2HIP.cpp index 920b310da8..7536c5c0f7 100644 --- a/hipamd/hipify-clang/src/CUDA2HIP.cpp +++ b/hipamd/hipify-clang/src/CUDA2HIP.cpp @@ -48,9 +48,9 @@ const std::map CUDA_INCLUDE_MAP{ {"curand_lognormal.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_mrg32k3a.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_mtgp32.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, - {"curand_mtgp32_host.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, + {"curand_mtgp32_host.h", {"hiprand_mtgp32_host.h", "", CONV_INCLUDE, API_RAND}}, {"curand_mtgp32_kernel.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, - {"curand_mtgp32dc_p_11213.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, + {"curand_mtgp32dc_p_11213.h", {"rocrand_mtgp32_11213.h", "", CONV_INCLUDE, API_RAND}}, {"curand_normal.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_normal_static.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, {"curand_philox4x32_x.h", {"hiprand_kernel.h", "", CONV_INCLUDE, API_RAND}}, diff --git a/hipamd/hipify-clang/src/CUDA2HIP_RAND_API_functions.cpp b/hipamd/hipify-clang/src/CUDA2HIP_RAND_API_functions.cpp index 65985d6f5c..525aa1e5f3 100644 --- a/hipamd/hipify-clang/src/CUDA2HIP_RAND_API_functions.cpp +++ b/hipamd/hipify-clang/src/CUDA2HIP_RAND_API_functions.cpp @@ -51,6 +51,8 @@ const std::map CUDA_RAND_FUNCTION_MAP{ {"curandSetPseudoRandomGeneratorSeed", {"hiprandSetPseudoRandomGeneratorSeed", "", CONV_LIB_FUNC, API_RAND}}, {"curandSetQuasiRandomGeneratorDimensions", {"hiprandSetQuasiRandomGeneratorDimensions", "", CONV_LIB_FUNC, API_RAND}}, {"curandSetStream", {"hiprandSetStream", "", CONV_LIB_FUNC, API_RAND}}, + {"curandMakeMTGP32Constants", {"hiprandMakeMTGP32Constants", "", CONV_LIB_FUNC, API_RAND}}, + {"curandMakeMTGP32KernelState", {"hiprandMakeMTGP32KernelState", "", CONV_LIB_FUNC, API_RAND}}, // RAND Device functions {"curand", {"hiprand", "", CONV_LIB_DEVICE_FUNC, API_RAND}}, diff --git a/hipamd/hipify-clang/src/HipifyAction.cpp b/hipamd/hipify-clang/src/HipifyAction.cpp index f4d68cd820..baa66e064c 100644 --- a/hipamd/hipify-clang/src/HipifyAction.cpp +++ b/hipamd/hipify-clang/src/HipifyAction.cpp @@ -191,8 +191,10 @@ bool HipifyAction::Exclude(const hipCounter & hipToken) { case CONV_INCLUDE: switch (hipToken.apiType) { case API_RAND: - if (insertedRAND_kernelHeader) { return true; } - insertedRAND_kernelHeader = true; + if (hipToken.hipName == "hiprand_kernel.h") { + if (insertedRAND_kernelHeader) { return true; } + insertedRAND_kernelHeader = true; + } return false; default: return false; diff --git a/hipamd/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp b/hipamd/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp index 9cb1a72420..b1cdec702a 100644 --- a/hipamd/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp +++ b/hipamd/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp @@ -37,9 +37,9 @@ #include // CHECK: #include #include -// CHECK-NOT: #include -// CHECK-NOT: #include +// CHECK: #include #include +// CHECK: #include #include // CHECK: if ((x) != hipSuccess) { @@ -201,11 +201,9 @@ struct runner CUDA_CALL(cudaMalloc((void **)&states, states_size * sizeof(curandStateMtgp32_t))); // CHECK: CUDA_CALL(hipMalloc((void **)&d_param, sizeof(mtgp32_kernel_params))); CUDA_CALL(cudaMalloc((void **)&d_param, sizeof(mtgp32_kernel_params))); - // curandMakeMTGP32Constants is yet unsupported by HIP - // CHECK-NOT: CURAND_CALL(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, d_param)); + // CHECK: CURAND_CALL(hiprandMakeMTGP32Constants(mtgp32dc_params_fast_11213, d_param)); CURAND_CALL(curandMakeMTGP32Constants(mtgp32dc_params_fast_11213, d_param)); - // curandMakeMTGP32KernelState is yet unsupported by HIP - // CHECK-NOT: CURAND_CALL(hiprandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, d_param, states_size, seed)); + // CHECK: CURAND_CALL(hiprandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, d_param, states_size, seed)); CURAND_CALL(curandMakeMTGP32KernelState(states, mtgp32dc_params_fast_11213, d_param, states_size, seed)); }