AI/rocm-systems
2
0
Vork 0
Code Kwesties Pull-aanvragen Actions Paketten Projecten Publicaties Wiki Activiteit

[HIPIFY][RAND] revise

Bladeren bron 
+ add missing functions
+ add minimum rocRAND support
+ updated CURAND_API_supported_by_HIP.md
This commit is contained in:
Evgeny Mankov
2019-02-08 17:54:28 +03:00
bovenliggende abb20bd382
commit d1ea9d9d97
5 gewijzigde bestanden met toevoegingen van 14 en 10 verwijderingen
Download Patch-bestand Download Diff-bestand Expand all files Collapse all files
hipamd/docs/markdown/CURAND_API_supported_by_HIP.md
+2
Bestand weergeven
@@ -134,6 +134,8 @@
| `curandSetPseudoRandomGeneratorSeed` | `hiprandSetPseudoRandomGeneratorSeed` |
| `curandSetQuasiRandomGeneratorDimensions` | `hiprandSetQuasiRandomGeneratorDimensions` |
| `curandSetStream` | `hiprandSetStream` |
| `curandMakeMTGP32Constants` | `hiprandMakeMTGP32Constants` |
| `curandMakeMTGP32KernelState` | `hiprandMakeMTGP32KernelState` |
## **3. Device API Functions**
hipamd/hipify-clang/src/CUDA2HIP.cpp
+2 -2
Bestand weergeven
@@ -48,9 +48,9 @@ const std::map <llvm::StringRef, hipCounter> 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}},
hipamd/hipify-clang/src/CUDA2HIP_RAND_API_functions.cpp
+2
Bestand weergeven
@@ -51,6 +51,8 @@ const std::map<llvm::StringRef, hipCounter> 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}},
hipamd/hipify-clang/src/HipifyAction.cpp
+4 -2
Bestand weergeven
@@ -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;
hipamd/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp
+4 -6
Bestand weergeven
@@ -37,9 +37,9 @@
#include <curand.h>
// CHECK: #include <hiprand_kernel.h>
#include <curand_kernel.h>
// CHECK-NOT: #include <curand_mtgp32_host.h>
// CHECK-NOT: #include <curand_mtgp32dc_p_11213.h>
// CHECK: #include <hiprand_mtgp32_host.h>
#include <curand_mtgp32_host.h>
// CHECK: #include <rocrand_mtgp32_11213.h>
#include <curand_mtgp32dc_p_11213.h>
// CHECK: if ((x) != hipSuccess) {
@@ -201,11 +201,9 @@ struct runner<curandStateMtgp32_t>
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));
}