Skip to content

rocFFT integration #139

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 14 commits into from
Aug 24, 2018
Merged

rocFFT integration #139

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
14 commits
Select commit Hold shift + click to select a range
84a985e
First integration patch to use rocFFT.
iotamudelta Aug 20, 2018
5e710fc
We don't have a license error. We are open source. :-)
iotamudelta Aug 20, 2018
c40247a
We do not have a comparable hip function (yet?) - ifdef to get wire it
iotamudelta Aug 20, 2018
777c095
Correct wrong syntax
iotamudelta Aug 20, 2018
1322aa1
Work around (not elegantly) the API not guaranteeing const for the
iotamudelta Aug 20, 2018
6098f41
Correctly use the API, consider that rocFFT needs ints, not long longs.
iotamudelta Aug 20, 2018
a02c34f
Work around missing hipfftXtExec API.
iotamudelta Aug 21, 2018
3bd1512
Merge branch 'master' into rocfft_integration
iotamudelta Aug 21, 2018
8f8b653
Fix ifdef
iotamudelta Aug 21, 2018
b67e1ee
Remove &.
iotamudelta Aug 21, 2018
103e534
As per @Jorghi12 request remove this.
iotamudelta Aug 21, 2018
8709047
Merge remote-tracking branch 'rocm_upstream/master'
iotamudelta Aug 23, 2018
1f0ab4b
Merge branch 'master' into rocfft_integration
iotamudelta Aug 24, 2018
6b85570
Merge branch 'master' into rocfft_integration
iotamudelta Aug 24, 2018
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion aten/src/ATen/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -247,7 +247,7 @@ IF(USE_CUDA AND NOT USE_ROCM)
ENDIF()

IF(USE_ROCM)
### Link in the ROCm libraries BLAS / RNG.
### Link in the ROCm libraries BLAS / RNG .
FIND_LIBRARY(ROCBLAS_LIBRARY rocblas HINTS ${ROCBLAS_PATH}/lib)
FIND_LIBRARY(HIPRAND_LIBRARY hiprand HINTS ${HIPRAND_PATH}/lib)

Expand Down
64 changes: 63 additions & 1 deletion aten/src/ATen/native/cuda/CuFFTPlanCache.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -149,7 +149,11 @@ class CuFFTConfig {
// TODO: Figure out why windows fails to compile
// at::optional<std::vector<long long int>> inembed_opt = at::nullopt;
// Then move the following to a helper function.
#ifdef __HIP_PLATFORM_HCC__
std::vector<int> inembed(signal_ndim);
#else
std::vector<long long int> inembed(signal_ndim);
#endif
if (!clone_input) {
auto istrides = input.strides();
auto last_istride = istrides[signal_ndim];
Expand Down Expand Up @@ -192,6 +196,37 @@ class CuFFTConfig {
inembed.begin()); // begin of output
}

#ifdef __HIP_PLATFORM_HCC__

hipfftType exec_type;
if (input.type().scalarType() == ScalarType::Float) {
if (complex_input && complex_output) {
exec_type = HIPFFT_C2C;
} else if (complex_input && !complex_output) {
exec_type = HIPFFT_C2R;
} else if (!complex_input && complex_output) {
exec_type = HIPFFT_R2C;
} else {
throw std::runtime_error("hipFFT doesn't support r2r (float)");
}
} else if (input.type().scalarType() == ScalarType::Double) {
if (complex_input && complex_output) {
exec_type = HIPFFT_Z2Z;
} else if (complex_input && !complex_output) {
exec_type = HIPFFT_Z2D;
} else if (!complex_input && complex_output) {
exec_type = HIPFFT_D2Z;
} else {
throw std::runtime_error("hipFFT doesn't support r2r (double)");
}
} else {
std::ostringstream ss;
ss << "hipFFT doesn't support tensor of type: "
<< at::toString(input.type().scalarType());
throw std::runtime_error(ss.str());
}

#else
cudaDataType itype, otype, exec_type;
if (input.type().scalarType() == ScalarType::Float) {
itype = complex_input ? CUDA_C_32F : CUDA_R_32F;
Expand All @@ -211,6 +246,7 @@ class CuFFTConfig {
<< at::toString(input.type().scalarType());
throw std::runtime_error(ss.str());
}
#endif

// create plan
auto raw_plan_ptr = new cufftHandle();
Expand All @@ -229,10 +265,18 @@ class CuFFTConfig {
// by assuming base_istride = base_ostride = 1.
//
// See NOTE [ cuFFT Embedded Strides ] in native/cuda/SpectralOps.cu.
#ifdef __HIP_PLATFORM_HCC__
int sizes = *signal_sizes.data();
Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Don't really need to create a new variable. Can simply pass signal_sizes.data().

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actually we do. signal_sizes.data() returns a long long* and rocFFT needs a int*

CUFFT_CHECK(hipfftMakePlanMany(plan(), signal_ndim, &sizes,
/* inembed */ nullptr, /* base_istride */ 1, /* idist */ 1,
/* onembed */ nullptr, /* base_ostride */ 1, /* odist */ 1,
exec_type, batch, &ws_size_t));
#else
CUFFT_CHECK(cufftXtMakePlanMany(plan(), signal_ndim, signal_sizes.data(),
/* inembed */ nullptr, /* base_istride */ 1, /* idist */ 1, itype,
/* onembed */ nullptr, /* base_ostride */ 1, /* odist */ 1, otype,
batch, &ws_size_t, exec_type));
#endif
} else {
// set idist (stride at batch dim)
// set base_istride (stride at innermost dim of signal)
Expand All @@ -254,6 +298,19 @@ class CuFFTConfig {
}

// set odist, onembed, base_ostride
#ifdef __HIP_PLATFORM_HCC__
int odist = at::prod_intlist(output_sizes.slice(1, signal_ndim));
std::vector<int> onembed(output_sizes.data() + 1, output_sizes.data() + signal_ndim + 1);
int base_ostride = 1;

int sizes = *signal_sizes.data();
int istride = base_istride;
int iidist = idist;
CUFFT_CHECK(hipfftMakePlanMany(plan(), signal_ndim, &sizes,
Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Don't really need to create a new variable. Can simply pass signal_sizes.data().

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actually we do. signal_sizes.data() returns a long long* and rocFFT needs a int*. Same goes for the other introduces variables. Our libraries team is aware that we need long long* API.

inembed.data(), istride, iidist,
onembed.data(), base_ostride, odist,
exec_type, batch, &ws_size_t));
#else
long long int odist = at::prod_intlist(output_sizes.slice(1, signal_ndim));
std::vector<long long int> onembed(output_sizes.data() + 1, output_sizes.data() + signal_ndim + 1);
long long int base_ostride = 1;
Expand All @@ -262,11 +319,16 @@ class CuFFTConfig {
inembed.data(), base_istride, idist, itype,
onembed.data(), base_ostride, odist, otype,
batch, &ws_size_t, exec_type));
}
#endif
}
ws_size = static_cast<int64_t>(ws_size_t);
}

#ifdef __HIP_PLATFORM_HCC__
cufftHandle &plan() const { return *plan_ptr.get(); }
Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why remove the constness here? The signatures for hipfftCreate & cufftCreate are identical.

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Because the signatures for MakeMany are not.

#else
const cufftHandle &plan() const { return *plan_ptr.get(); }
#endif

bool should_clone_input() const { return clone_input; }

Expand Down
2 changes: 2 additions & 0 deletions aten/src/ATen/native/cuda/CuFFTUtils.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,8 +49,10 @@ static inline std::string _cudaGetErrorEnum(cufftResult error)
return "CUFFT_NO_WORKSPACE";
case CUFFT_NOT_IMPLEMENTED:
return "CUFFT_NOT_IMPLEMENTED";
#ifndef __HIP_PLATFORM_HCC__
case CUFFT_LICENSE_ERROR:
return "CUFFT_LICENSE_ERROR";
#endif
case CUFFT_NOT_SUPPORTED:
return "CUFFT_NOT_SUPPORTED";
default:
Expand Down
37 changes: 37 additions & 0 deletions aten/src/ATen/native/cuda/SpectralOps.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -190,8 +190,45 @@ static inline Tensor _run_cufft(
CUFFT_CHECK(cufftSetWorkArea(plan, ws.data_ptr()));

// run
#ifdef __HIP_PLATFORM_HCC__
Copy link

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is fine for now but would be good to file an issue under the rocFFT to further extend their API for a cufftXtExec-esque call.

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Already done....

if (input.type().scalarType() == ScalarType::Float) {
if (complex_input && complex_output) {
CUFFT_CHECK(hipfftExecC2C(plan, static_cast<hipfftComplex*>(input.data_ptr()),
static_cast<hipfftComplex*>(output.data_ptr()),
inverse ? HIPFFT_BACKWARD : HIPFFT_FORWARD));
} else if (complex_input && !complex_output) {
CUFFT_CHECK(hipfftExecC2R(plan, static_cast<hipfftComplex*>(input.data_ptr()),
static_cast<hipfftReal*>(output.data_ptr())));
} else if (!complex_input && complex_output) {
CUFFT_CHECK(hipfftExecR2C(plan, static_cast<hipfftReal*>(input.data_ptr()),
static_cast<hipfftComplex*>(output.data_ptr())));
} else {
throw std::runtime_error("hipFFT doesn't support r2r (float)");
}
} else if (input.type().scalarType() == ScalarType::Double) {
if (complex_input && complex_output) {
CUFFT_CHECK(hipfftExecZ2Z(plan, static_cast<hipfftDoubleComplex*>(input.data_ptr()),
static_cast<hipfftDoubleComplex*>(output.data_ptr()),
inverse ? HIPFFT_BACKWARD : HIPFFT_FORWARD));
} else if (complex_input && !complex_output) {
CUFFT_CHECK(hipfftExecZ2D(plan, static_cast<hipfftDoubleComplex*>(input.data_ptr()),
static_cast<hipfftDoubleReal*>(output.data_ptr())));
} else if (!complex_input && complex_output) {
CUFFT_CHECK(hipfftExecD2Z(plan, static_cast<hipfftDoubleReal*>(input.data_ptr()),
static_cast<hipfftDoubleComplex*>(output.data_ptr())));
} else {
throw std::runtime_error("hipFFT doesn't support r2r (double)");
}
} else {
std::ostringstream ss;
ss << "hipFFT doesn't support tensor of type: "
<< at::toString(input.type().scalarType());
throw std::runtime_error(ss.str());
}
#else
CUFFT_CHECK(cufftXtExec(plan, input.data_ptr(), output.data_ptr(),
inverse ? CUFFT_INVERSE : CUFFT_FORWARD));
#endif

// rescale if needed by normalized flag or inverse transform
auto size_last_signal_dim = checked_signal_sizes[signal_ndim - 1];
Expand Down
1 change: 1 addition & 0 deletions cmake/Dependencies.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -562,6 +562,7 @@ endif()
if(USE_ROCM)
include_directories(SYSTEM ${HIP_PATH}/include)
include_directories(SYSTEM ${ROCBLAS_PATH}/include)
include_directories(SYSTEM ${ROCFFT_PATH}/include)
include_directories(SYSTEM ${HIPSPARSE_PATH}/include)
include_directories(SYSTEM ${HIPRAND_PATH}/include)
include_directories(SYSTEM ${ROCRAND_PATH}/include)
Expand Down
9 changes: 9 additions & 0 deletions cmake/public/LoadHIP.cmake
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -38,6 +38,13 @@ ELSE()
SET(ROCBLAS_PATH $ENV{ROCBLAS_PATH})
ENDIF()

# ROCFFT_PATH
IF(NOT DEFINED ENV{ROCFFT_PATH})
SET(ROCBLAS_PATH ${ROCM_PATH}/rocfft)
ELSE()
SET(ROCFFT_PATH $ENV{ROCFFT_PATH})
ENDIF()

# HIPSPARSE_PATH
IF(NOT DEFINED ENV{HIPSPARSE_PATH})
SET(HIPSPARSE_PATH ${ROCM_PATH}/hcsparse)
Expand Down Expand Up @@ -106,11 +113,13 @@ IF(HIP_FOUND)
set(rocblas_DIR ${ROCBLAS_PATH}/lib/cmake/rocblas)
set(miopen_DIR ${MIOPEN_PATH}/lib/cmake/miopen)
set(rocblas_DIR ${ROCBLAS_PATH}/lib/cmake/rocblas)
set(rocfft_DIR ${ROCFFT_PATH}/lib/cmake/rocfft)
set(hipsparse_DIR ${HIPSPARSE_PATH}/lib/cmake/hipsparse)

find_package(rocrand REQUIRED)
find_package(hiprand REQUIRED)
find_package(rocblas REQUIRED)
find_package(rocfft REQUIRED)
find_package(miopen REQUIRED)
#find_package(hipsparse REQUIRED)

Expand Down
2 changes: 2 additions & 0 deletions setup.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -920,6 +920,7 @@ def run(self):
rocm_include_path = '/opt/rocm/include'
hcc_include_path = '/opt/rocm/hcc/include'
rocblas_include_path = '/opt/rocm/rocblas/include'
rocfft_include_path = '/opt/rocm/rocfft/include'
hipsparse_include_path = '/opt/rocm/hcsparse/include'
hiprand_include_path = '/opt/rocm/hiprand/include'
rocrand_include_path = '/opt/rocm/rocrand/include'
Expand All @@ -928,6 +929,7 @@ def run(self):
include_dirs.append(rocm_include_path)
include_dirs.append(hcc_include_path)
include_dirs.append(rocblas_include_path)
include_dirs.append(rocfft_include_path)
include_dirs.append(hipsparse_include_path)
include_dirs.append(hiprand_include_path)
include_dirs.append(rocrand_include_path)
Expand Down
24 changes: 0 additions & 24 deletions tools/amd_build/disabled_features.yaml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -97,27 +97,6 @@
"struct mtgp32_kernel_params": "mtgp32_kernel_params"
}
},
{
"path": "aten/src/ATen/native/cuda/CuFFTUtils.h",
"s_constants": {
"#include <cufft.h>": "",
"#include <cufftXt.h>": ""
}
},
{
"path": "aten/src/ATen/native/cuda/CuFFTPlanCache.h",
"s_constants": {
"#include <cufft.h>": "",
"#include <cufftXt.h>": ""
}
},
{
"path": "aten/src/ATen/native/cuda/SpectralOps.cu",
"s_constants": {
"#include <cufft.h>": "",
"#include <cufftXt.h>": ""
}
},
{
"path": "aten/src/ATen/native/cuda/RoiPooling.cu",
"s_constants": {
Expand All @@ -142,9 +121,6 @@
}
],
"disabled_modules": [
"aten/src/ATen/native/cuda/CuFFTUtils.h",
"aten/src/ATen/native/cuda/CuFFTPlanCache.h",
"aten/src/ATen/native/cuda/SpectralOps.cu",
],
"disabled_functions": [
{
Expand Down
3 changes: 2 additions & 1 deletion tools/amd_build/pyHIPIFY/constants.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,7 +50,8 @@
API_SPARSE = 40
API_RAND = 41
API_LAST = 42
API_FFT = 43

HIP_UNSUPPORTED = 43
API_PYTORCH = 1337
API_CAFFE2 = 1338
API_CAFFE2 = 1338
76 changes: 73 additions & 3 deletions tools/amd_build/pyHIPIFY/cuda_to_hip_mappings.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -272,8 +272,8 @@
"curand_precalc.h": ("hiprand_kernel.h", CONV_INCLUDE, API_RAND),
"curand_uniform.h": ("hiprand_kernel.h", CONV_INCLUDE, API_RAND),
"cusparse.h": ("hipsparse.h", CONV_INCLUDE, API_RAND),
"#include <cufft.h>": ("", CONV_INCLUDE, API_RAND, HIP_UNSUPPORTED),
"#include <cufftXt.h>": ("", CONV_INCLUDE, API_RAND, HIP_UNSUPPORTED),
"cufft.h": ("hipfft.h", CONV_INCLUDE, API_BLAS),
"cufftXt.h": ("hipfft.h", CONV_INCLUDE, API_BLAS),
"#include <nvfunctional>": ("", CONV_INCLUDE, API_RAND, HIP_UNSUPPORTED),
}

Expand Down Expand Up @@ -2096,7 +2096,77 @@
"curand_poisson": ("hiprand_poisson", CONV_DEVICE_FUNC, API_RAND),
"curand_poisson4": ("hiprand_poisson4", CONV_DEVICE_FUNC, API_RAND),
"curand_Philox4x32_10": ("hiprand_Philox4x32_10", CONV_DEVICE_FUNC, API_RAND, HIP_UNSUPPORTED),
"mtgp32_kernel_params": ("mtgp32_kernel_params_t", CONV_MATH_FUNC, API_RAND)
"mtgp32_kernel_params": ("mtgp32_kernel_params_t", CONV_MATH_FUNC, API_RAND),
"CUFFT_FORWARD": ("HIPFFT_FORWARD", CONV_NUMERIC_LITERAL, API_BLAS),
"CUFFT_INVERSE": ("HIPFFT_BACKWARD", CONV_NUMERIC_LITERAL, API_BLAS),
"CUFFT_COMPATIBILITY_DEFAULT": ("HIPFFT_COMPATIBILITY_DEFAULT", CONV_NUMERIC_LITERAL, API_BLAS, HIP_UNSUPPORTED),
"cufftResult_t": ("hipfftResult_t", CONV_TYPE, API_FFT),
"cufftResult": ("hipfftResult", CONV_TYPE, API_FFT),
"CUFFT_SUCCESS": ("HIPFFT_SUCCESS", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_INVALID_PLAN": ("HIPFFT_INVALID_PLAN", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_ALLOC_FAILED": ("HIPFFT_ALLOC_FAILED", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_INVALID_TYPE": ("HIPFFT_INVALID_TYPE", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_INVALID_VALUE": ("HIPFFT_INVALID_VALUE", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_INTERNAL_ERROR": ("HIPFFT_INTERNAL_ERROR", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_EXEC_FAILED": ("HIPFFT_EXEC_FAILED", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_SETUP_FAILED": ("HIPFFT_SETUP_FAILED", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_INVALID_SIZE": ("HIPFFT_INVALID_SIZE", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_UNALIGNED_DATA": ("HIPFFT_UNALIGNED_DATA", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_INCOMPLETE_PARAMETER_LIST": ("HIPFFT_INCOMPLETE_PARAMETER_LIST", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_INVALID_DEVICE": ("HIPFFT_INVALID_DEVICE", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_PARSE_ERROR": ("HIPFFT_PARSE_ERROR", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_NO_WORKSPACE": ("HIPFFT_NO_WORKSPACE", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_NOT_IMPLEMENTED": ("HIPFFT_NOT_IMPLEMENTED", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_LICENSE_ERROR": ("HIPFFT_LICENSE_ERROR", CONV_NUMERIC_LITERAL, API_FFT, HIP_UNSUPPORTED),
"CUFFT_NOT_SUPPORTED": ("HIPFFT_NOT_SUPPORTED", CONV_NUMERIC_LITERAL, API_FFT),
"cufftType_t": ("hipfftType_t", CONV_TYPE, API_FFT),
"cufftType": ("hipfftType", CONV_TYPE, API_FFT),
"CUFFT_R2C": ("HIPFFT_R2C", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_C2R": ("HIPFFT_C2R", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_C2C": ("HIPFFT_C2C", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_D2Z": ("HIPFFT_D2Z", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_Z2D": ("HIPFFT_Z2D", CONV_NUMERIC_LITERAL, API_FFT),
"CUFFT_Z2Z": ("HIPFFT_Z2Z", CONV_NUMERIC_LITERAL, API_FFT),
"cufftCompatibility_t": ("hipfftCompatibility_t", CONV_TYPE, API_FFT, HIP_UNSUPPORTED),
"cufftCompatibility": ("hipfftCompatibility", CONV_TYPE, API_FFT, HIP_UNSUPPORTED),
"CUFFT_COMPATIBILITY_FFTW_PADDING": ("HIPFFT_COMPATIBILITY_FFTW_PADDING", CONV_NUMERIC_LITERAL, API_FFT, HIP_UNSUPPORTED),
"cufftReal": ("hipfftReal", CONV_TYPE, API_FFT),
"cufftDoubleReal": ("hipfftDoubleReal", CONV_TYPE, API_FFT),
"cufftComplex": ("hipfftComplex", CONV_TYPE, API_FFT),
"cufftDoubleComplex": ("hipfftDoubleComplex", CONV_TYPE, API_FFT),
"cufftHandle": ("hipfftHandle", CONV_TYPE, API_FFT),
"cufftPlan1d": ("hipfftPlan1d", CONV_MATH_FUNC, API_FFT),
"cufftPlan2d": ("hipfftPlan2d", CONV_MATH_FUNC, API_FFT),
"cufftPlan3d": ("hipfftPlan3d", CONV_MATH_FUNC, API_FFT),
"cufftPlanMany": ("hipfftPlanMany", CONV_MATH_FUNC, API_FFT),
"cufftMakePlan1d": ("hipfftMakePlan1d", CONV_MATH_FUNC, API_FFT),
"cufftMakePlan2d": ("hipfftMakePlan2d", CONV_MATH_FUNC, API_FFT),
"cufftMakePlan3d": ("hipfftMakePlan3d", CONV_MATH_FUNC, API_FFT),
"cufftMakePlanMany": ("hipfftMakePlanMany", CONV_MATH_FUNC, API_FFT),
"cufftMakePlanMany64": ("hipfftMakePlanMany64", CONV_MATH_FUNC, API_FFT),
"cufftGetSizeMany64": ("hipfftGetSizeMany64", CONV_MATH_FUNC, API_FFT),
"cufftEstimate1d": ("hipfftEstimate1d", CONV_MATH_FUNC, API_FFT),
"cufftEstimate2d": ("hipfftEstimate2d", CONV_MATH_FUNC, API_FFT),
"cufftEstimate3d": ("hipfftEstimate3d", CONV_MATH_FUNC, API_FFT),
"cufftEstimateMany": ("hipfftEstimateMany", CONV_MATH_FUNC, API_FFT),
"cufftCreate": ("hipfftCreate", CONV_MATH_FUNC, API_FFT),
"cufftGetSize1d": ("hipfftGetSize1d", CONV_MATH_FUNC, API_FFT),
"cufftGetSize2d": ("hipfftGetSize2d", CONV_MATH_FUNC, API_FFT),
"cufftGetSize3d": ("hipfftGetSize3d", CONV_MATH_FUNC, API_FFT),
"cufftGetSizeMany": ("hipfftGetSizeMany", CONV_MATH_FUNC, API_FFT),
"cufftGetSize": ("hipfftGetSize", CONV_MATH_FUNC, API_FFT),
"cufftSetWorkArea": ("hipfftSetWorkArea", CONV_MATH_FUNC, API_FFT),
"cufftSetAutoAllocation": ("hipfftSetAutoAllocation", CONV_MATH_FUNC, API_FFT),
"cufftExecC2C": ("hipfftExecC2C", CONV_MATH_FUNC, API_FFT),
"cufftExecR2C": ("hipfftExecR2C", CONV_MATH_FUNC, API_FFT),
"cufftExecC2R": ("hipfftExecC2R", CONV_MATH_FUNC, API_FFT),
"cufftExecZ2Z": ("hipfftExecZ2Z", CONV_MATH_FUNC, API_FFT),
"cufftExecD2Z": ("hipfftExecD2Z", CONV_MATH_FUNC, API_FFT),
"cufftExecZ2D": ("hipfftExecZ2D", CONV_MATH_FUNC, API_FFT),
"cufftSetStream": ("hipfftSetStream", CONV_MATH_FUNC, API_FFT),
"cufftDestroy": ("hipfftDestroy", CONV_MATH_FUNC, API_FFT),
"cufftGetVersion": ("hipfftGetVersion", CONV_MATH_FUNC, API_FFT),
"cufftGetProperty": ("hipfftGetProperty", CONV_MATH_FUNC, API_FFT, HIP_UNSUPPORTED),
}

CUDA_SPARSE_MAP = {
Expand Down