Merge pull request #624 from emankov/cuComplex
[HIPIFY][Complex] Add cuComplex support
[ROCm/hip commit: 3ddfbfa94a]
Этот коммит содержится в:
Evgeny Mankov
GitHub
@@ -0,0 +1,37 @@
|
||||
# cuComplex API supported by HIP
|
||||
|
||||
## **1. cuComplex Data types**
|
||||
|
||||
| **type** | **CUDA** | **HIP** |**HIP value** (if differs) |
|
||||
|-------------:|---------------------------------------------------------------|------------------------------------------------------------|---------------------------|
|
||||
| float2 |***`cuFloatComplex`*** |***`hipFloatComplex`*** | struct |
|
||||
| double2 |***`cuDoubleComplex`*** |***`hipDoubleComplex`*** | struct |
|
||||
| float2 |***`cuComplex`*** |***`hipComplex`*** | struct |
|
||||
|
||||
## **2. cuComplex API functions**
|
||||
|
||||
| **CUDA** | **HIP** |
|
||||
|-----------------------------------------------------------|-------------------------------------------------|
|
||||
|`cuCrealf` |`hipCrealf` |
|
||||
|`cuCimagf` |`hipCimagf` |
|
||||
|`make_cuFloatComplex` |`make_hipFloatComplex` |
|
||||
|`cuConjf` |`hipConjf` |
|
||||
|`cuCaddf` |`hipCaddf` |
|
||||
|`cuCsubf` |`hipCsubf` |
|
||||
|`cuCmulf` |`hipCmulf` |
|
||||
|`cuCdivf` |`hipCdivf` |
|
||||
|`cuCabsf` |`hipCabsf` |
|
||||
|`cuCreal` |`hipCreal` |
|
||||
|`cuCimag` |`hipCimag` |
|
||||
|`make_cuDoubleComplex` |`make_hipDoubleComplex` |
|
||||
|`cuConj` |`hipConj` |
|
||||
|`cuCadd` |`hipCadd` |
|
||||
|`cuCsub` |`hipCsub` |
|
||||
|`cuCmul` |`hipCmul` |
|
||||
|`cuCdiv` |`hipCdiv` |
|
||||
|`cuCabs` |`hipCabs` |
|
||||
|`make_cuComplex` |`make_hipComplex` |
|
||||
|`cuComplexFloatToDouble` |`hipComplexFloatToDouble` |
|
||||
|`cuComplexDoubleToFloat` |`hipComplexDoubleToFloat` |
|
||||
|`cuCfmaf` |`hipCfmaf` |
|
||||
|`cuCfma` |`hipCfma` |
|
||||
@@ -21,6 +21,7 @@
|
||||
|
||||
- [Runtime API](../docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md)
|
||||
- [Driver API](../docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md)
|
||||
- [cuComplex API](../docs/markdown/cuComplex_API_supported_by_HIP.md)
|
||||
- [cuBLAS](../docs/markdown/CUBLAS_API_supported_by_HIP.md)
|
||||
- [cuRAND](../docs/markdown/CURAND_API_supported_by_HIP.md)
|
||||
- [cuDNN](../docs/markdown/CUDNN_API_supported_by_HIP.md)
|
||||
|
||||
@@ -366,7 +366,6 @@ const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP{
|
||||
{"channel_descriptor.h", {"hip/channel_descriptor.h", CONV_INCLUDE, API_RUNTIME}},
|
||||
{"device_functions.h", {"hip/device_functions.h", CONV_INCLUDE, API_RUNTIME}},
|
||||
{"driver_types.h", {"hip/driver_types.h", CONV_INCLUDE, API_RUNTIME}},
|
||||
{"cuComplex.h", {"hip/hip_complex.h", CONV_INCLUDE, API_RUNTIME}},
|
||||
{"cuda_fp16.h", {"hip/hip_fp16.h", CONV_INCLUDE, API_RUNTIME}},
|
||||
{"cuda_texture_types.h", {"hip/hip_texture_types.h", CONV_INCLUDE, API_RUNTIME}},
|
||||
{"vector_types.h", {"hip/hip_vector_types.h", CONV_INCLUDE, API_RUNTIME}},
|
||||
@@ -397,12 +396,11 @@ const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP{
|
||||
// CUDNN includes
|
||||
{"cudnn.h", {"hipDNN.h", CONV_INCLUDE_CUDA_MAIN_H, API_DNN}},
|
||||
|
||||
// CUDNN includes
|
||||
// CUFFT includes
|
||||
{"cufft.h", {"hipfft.h", CONV_INCLUDE_CUDA_MAIN_H, API_FFT}},
|
||||
|
||||
// HIP includes
|
||||
// TODO: uncomment this when hip/cudacommon.h will be renamed to hip/hipcommon.h
|
||||
// {"cudacommon.h", {"hipcommon.h", CONV_INCLUDE, API_RUNTIME}},
|
||||
// cuComplex includes
|
||||
{"cuComplex.h", {"hip/hip_complex.h", CONV_INCLUDE_CUDA_MAIN_H, API_COMPLEX}},
|
||||
};
|
||||
|
||||
/// All other identifiers: function and macro names.
|
||||
@@ -1387,6 +1385,36 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
|
||||
{"cuGraphicsEGLRegisterImage", {"hipGraphicsEGLRegisterImage", CONV_EGL, API_DRIVER, HIP_UNSUPPORTED}}, // API_Runtime ANALOGUE (cudaGraphicsEGLRegisterImage)
|
||||
{"cuGraphicsResourceGetMappedEglFrame", {"hipGraphicsResourceGetMappedEglFrame", CONV_EGL, API_DRIVER, HIP_UNSUPPORTED}}, // API_Runtime ANALOGUE (cudaGraphicsResourceGetMappedEglFrame)
|
||||
|
||||
|
||||
////////////////////////////// cuComplex API //////////////////////////////
|
||||
{"cuFloatComplex", {"hipFloatComplex", CONV_TYPE, API_COMPLEX}},
|
||||
{"cuDoubleComplex", {"hipDoubleComplex", CONV_TYPE, API_COMPLEX}},
|
||||
{"cuComplex", {"hipComplex", CONV_TYPE, API_COMPLEX}},
|
||||
|
||||
{"cuCrealf", {"hipCrealf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCimagf", {"hipCimagf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"make_cuFloatComplex", {"make_hipFloatComplex", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuConjf", {"hipConjf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCaddf", {"hipCaddf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCsubf", {"hipCsubf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCmulf", {"hipCmulf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCdivf", {"hipCdivf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCabsf", {"hipCabsf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCreal", {"hipCreal", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCimag", {"hipCimag", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"make_cuDoubleComplex", {"make_hipDoubleComplex", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuConj", {"hipConj", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCadd", {"hipCadd", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCsub", {"hipCsub", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCmul", {"hipCmul", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCdiv", {"hipCdiv", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCabs", {"hipCabs", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"make_cuComplex", {"make_hipComplex", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuComplexFloatToDouble", {"hipComplexFloatToDouble", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuComplexDoubleToFloat", {"hipComplexDoubleToFloat", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCfmaf", {"hipCfmaf", CONV_COMPLEX, API_COMPLEX}},
|
||||
{"cuCfma", {"hipCfma", CONV_COMPLEX, API_COMPLEX}},
|
||||
|
||||
/////////////////////////////// CUDA RT API ///////////////////////////////
|
||||
// Data types
|
||||
{"cudaDataType_t", {"hipDataType_t", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}},
|
||||
|
||||
@@ -168,6 +168,10 @@ bool HipifyAction::Exclude(const hipCounter & hipToken) {
|
||||
if (insertedFFTHeader) { return true; }
|
||||
insertedFFTHeader = true;
|
||||
return false;
|
||||
case API_COMPLEX:
|
||||
if (insertedComplexHeader) { return true; }
|
||||
insertedComplexHeader = true;
|
||||
return false;
|
||||
default:
|
||||
return false;
|
||||
}
|
||||
|
||||
@@ -29,6 +29,7 @@ private:
|
||||
bool insertedRAND_kernelHeader = false;
|
||||
bool insertedDNNHeader = false;
|
||||
bool insertedFFTHeader = false;
|
||||
bool insertedComplexHeader = false;
|
||||
bool firstHeader = false;
|
||||
bool pragmaOnce = false;
|
||||
clang::SourceLocation firstHeaderLoc;
|
||||
|
||||
@@ -8,13 +8,13 @@ const char *counterNames[NUM_CONV_TYPES] = {
|
||||
"version", "init", "device", "mem", "kern", "coord_func", "math_func", "device_func",
|
||||
"special_func", "stream", "event", "occupancy", "ctx", "peer", "module",
|
||||
"cache", "exec", "err", "def", "tex", "gl", "graphics",
|
||||
"surface", "jit", "d3d9", "d3d10", "d3d11", "vdpau", "egl",
|
||||
"surface", "jit", "d3d9", "d3d10", "d3d11", "vdpau", "egl", "complex",
|
||||
"thread", "other", "include", "include_cuda_main_header", "type", "literal",
|
||||
"numeric_literal"
|
||||
};
|
||||
|
||||
const char *apiNames[NUM_API_TYPES] = {
|
||||
"CUDA Driver API", "CUDA RT API", "CUBLAS API", "CURAND API", "CUDNN API", "CUFFT API"
|
||||
"CUDA Driver API", "CUDA RT API", "CUBLAS API", "CURAND API", "CUDNN API", "CUFFT API", "cuComplex API"
|
||||
};
|
||||
|
||||
namespace {
|
||||
|
||||
@@ -40,6 +40,7 @@ enum ConvTypes {
|
||||
CONV_D3D11,
|
||||
CONV_VDPAU,
|
||||
CONV_EGL,
|
||||
CONV_COMPLEX,
|
||||
CONV_THREAD,
|
||||
CONV_OTHER,
|
||||
CONV_INCLUDE,
|
||||
@@ -58,6 +59,7 @@ enum ApiTypes {
|
||||
API_RAND,
|
||||
API_DNN,
|
||||
API_FFT,
|
||||
API_COMPLEX,
|
||||
API_LAST
|
||||
};
|
||||
constexpr int NUM_API_TYPES = (int) ApiTypes::API_LAST;
|
||||
|
||||
@@ -0,0 +1,58 @@
|
||||
// RUN: %run_test hipify "%s" "%t" %cuda_args
|
||||
|
||||
// CHECK: #include <hip/hip_runtime.h>
|
||||
// CHECK: #include "hip/hip_complex.h"
|
||||
#include "cuComplex.h"
|
||||
|
||||
#define TYPEFLOAT
|
||||
#define DIMX 100
|
||||
#define DIMY 40
|
||||
#define moveX 2
|
||||
#define moveY 1
|
||||
|
||||
#define MAXITERATIONS 10
|
||||
|
||||
#ifdef TYPEFLOAT
|
||||
#define TYPE float
|
||||
// CHECK: #define cTYPE hipFloatComplex
|
||||
#define cTYPE cuFloatComplex
|
||||
// CHECK: #define cMakecuComplex(re,i) make_hipFloatComplex(re,i)
|
||||
#define cMakecuComplex(re,i) make_cuFloatComplex(re,i)
|
||||
#endif
|
||||
#ifdef TYPEDOUBLE
|
||||
// CHECK: #define TYPE hipDoubleComplex
|
||||
#define TYPE cuDoubleComplex
|
||||
// CHECK: #define cMakecuComplex(re,i) make_hipDoubleComplex(re,i)
|
||||
#define cMakecuComplex(re,i) make_cuDoubleComplex(re,i)
|
||||
#endif
|
||||
|
||||
__device__ cTYPE juliaFunctor(cTYPE p, cTYPE c) {
|
||||
// CHECK: return hipCaddf(hipCmulf(p, p), c);
|
||||
return cuCaddf(cuCmulf(p, p), c);
|
||||
}
|
||||
|
||||
__device__ cTYPE convertToComplex(int x, int y, float zoom) {
|
||||
TYPE jx = 1.5 * (x - DIMX / 2) / (0.5 * zoom * DIMX) + moveX;
|
||||
TYPE jy = (y - DIMY / 2) / (0.5 * zoom * DIMY) + moveY;
|
||||
return cMakecuComplex(jx, jy);
|
||||
}
|
||||
|
||||
__device__ int evolveComplexPoint(cTYPE p, cTYPE c) {
|
||||
int it = 1;
|
||||
// CHECK: while (it <= MAXITERATIONS && hipCabsf(p) <= 4) {
|
||||
while (it <= MAXITERATIONS && cuCabsf(p) <= 4) {
|
||||
p = juliaFunctor(p, c);
|
||||
it++;
|
||||
}
|
||||
return it;
|
||||
}
|
||||
|
||||
__global__ void computeJulia(int* data, cTYPE c, float zoom) {
|
||||
int i = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
int j = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
|
||||
if (i<DIMX && j<DIMY) {
|
||||
cTYPE p = convertToComplex(i, j, zoom);
|
||||
data[i*DIMY + j] = evolveComplexPoint(p, c);
|
||||
}
|
||||
}
|
||||
Ссылка в новой задаче
Block a user