Skip to content

Commit

Permalink
[HIPIFY][rocRAND][feature] Support for cuRAND -> rocRAND hipificati…
Browse files Browse the repository at this point in the history
…on - Step 2 - Headers

+ Enumerate all `rocRAND` headers
+ Provided a corresponding test `headers_test_09_rocrand.cu`

[ToDo]
+ [perl] Fix roc headers substitution under the `--roc` option
  • Loading branch information
emankov committed Aug 16, 2024
1 parent b584e37 commit f70c820
Show file tree
Hide file tree
Showing 8 changed files with 167 additions and 32 deletions.
2 changes: 1 addition & 1 deletion bin/hipify-perl
Original file line number Diff line number Diff line change
Expand Up @@ -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");
Expand Down
34 changes: 17 additions & 17 deletions src/CUDA2HIP.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -48,23 +48,23 @@ const std::map <llvm::StringRef, hipCounter> 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
Expand Down
42 changes: 31 additions & 11 deletions src/HipifyAction.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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]";
Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -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:
Expand Down
3 changes: 3 additions & 0 deletions src/HipifyAction.h
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down
2 changes: 1 addition & 1 deletion tests/unit_tests/headers/headers_test_09.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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 <string>

Expand Down
2 changes: 1 addition & 1 deletion tests/unit_tests/headers/headers_test_09_12000.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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 <string>

Expand Down
112 changes: 112 additions & 0 deletions tests/unit_tests/headers/headers_test_09_rocrand.cu
Original file line number Diff line number Diff line change
@@ -0,0 +1,112 @@
// RUN: %run_test hipify "%s" "%t" %hipify_args 1 --roc %clang_args

// CHECK: #include <hip/hip_runtime.h>
// CHECK: #include <memory>

// CHECK-NOT: #include <cuda_runtime.h>
// CHECK-NOT: #include <hip/hip_runtime.h>

// 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 <iostream>

// CHECK: #include <stdio.h>

// CHECK: #include "rocrand/rocrand.h"
// CHECK: #include "rocrand/rocrand_kernel.h"

// CHECK: #include <algorithm>

// 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 <string>

// CHECK: #include "hipfft/hipfft.h"
// CHECK: #include "rocsparse.h"

#include <cuda.h>
// CHECK-NOT: #include <hip/hip_runtime.h>

#include <memory>

#include <cuda_runtime.h>
// CHECK-NOT: #include <hip/hip_runtime.h>

#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 <iostream>

#include <stdio.h>

#include "curand.h"
#include "curand_kernel.h"

#include <algorithm>

#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 <string>

#include "cufft.h"

#include "cusparse.h"
Original file line number Diff line number Diff line change
Expand Up @@ -39,7 +39,7 @@
#include <curand_kernel.h>
// CHECK: #include <hiprand/hiprand_mtgp32_host.h>
#include <curand_mtgp32_host.h>
// CHECK: #include <rocrand_mtgp32_11213.h>
// CHECK: #include <rocrand/rocrand_mtgp32_11213.h>
#include <curand_mtgp32dc_p_11213.h>

// CHECK: if ((x) != hipSuccess) {
Expand Down

0 comments on commit f70c820

Please sign in to comment.