From 37d047b454a07a72598c87ce25ab3788cd152d69 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 8 Aug 2018 18:34:57 +0300 Subject: [PATCH] [HIPIFY][Complex] Add cuComplex support + Add API_COMPLEX support (data types and functions) + Add cuComplex_API_supported_by_HIP.md + Add cuComplex_Julia.cu test + Update README.md [ROCm/hip commit: 368977f75bfdfb478dab169a635148e0fbc372b9] --- .../cuComplex_API_supported_by_HIP.md | 37 ++++++++++++ projects/hip/hipify-clang/README.md | 1 + projects/hip/hipify-clang/src/CUDA2HipMap.cpp | 38 ++++++++++-- .../hip/hipify-clang/src/HipifyAction.cpp | 4 ++ projects/hip/hipify-clang/src/HipifyAction.h | 1 + projects/hip/hipify-clang/src/Statistics.cpp | 4 +- projects/hip/hipify-clang/src/Statistics.h | 2 + .../hipify-clang/cuComplex/cuComplex_Julia.cu | 58 +++++++++++++++++++ 8 files changed, 138 insertions(+), 7 deletions(-) create mode 100644 projects/hip/docs/markdown/cuComplex_API_supported_by_HIP.md create mode 100644 projects/hip/tests/hipify-clang/cuComplex/cuComplex_Julia.cu diff --git a/projects/hip/docs/markdown/cuComplex_API_supported_by_HIP.md b/projects/hip/docs/markdown/cuComplex_API_supported_by_HIP.md new file mode 100644 index 0000000000..621973ba61 --- /dev/null +++ b/projects/hip/docs/markdown/cuComplex_API_supported_by_HIP.md @@ -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` | diff --git a/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md index d5a825033d..7f55dea800 100644 --- a/projects/hip/hipify-clang/README.md +++ b/projects/hip/hipify-clang/README.md @@ -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) diff --git a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp index ca53738abb..4bd7263bad 100644 --- a/projects/hip/hipify-clang/src/CUDA2HipMap.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HipMap.cpp @@ -366,7 +366,6 @@ const std::map 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 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 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}}, diff --git a/projects/hip/hipify-clang/src/HipifyAction.cpp b/projects/hip/hipify-clang/src/HipifyAction.cpp index 20bc9dcc2c..a9bd1aa085 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.cpp +++ b/projects/hip/hipify-clang/src/HipifyAction.cpp @@ -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; } diff --git a/projects/hip/hipify-clang/src/HipifyAction.h b/projects/hip/hipify-clang/src/HipifyAction.h index ad987c921e..1262142cfc 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.h +++ b/projects/hip/hipify-clang/src/HipifyAction.h @@ -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; diff --git a/projects/hip/hipify-clang/src/Statistics.cpp b/projects/hip/hipify-clang/src/Statistics.cpp index 595f243857..2115c567a5 100644 --- a/projects/hip/hipify-clang/src/Statistics.cpp +++ b/projects/hip/hipify-clang/src/Statistics.cpp @@ -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 { diff --git a/projects/hip/hipify-clang/src/Statistics.h b/projects/hip/hipify-clang/src/Statistics.h index 00ecce3eda..1ded45f0e4 100644 --- a/projects/hip/hipify-clang/src/Statistics.h +++ b/projects/hip/hipify-clang/src/Statistics.h @@ -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; diff --git a/projects/hip/tests/hipify-clang/cuComplex/cuComplex_Julia.cu b/projects/hip/tests/hipify-clang/cuComplex/cuComplex_Julia.cu new file mode 100644 index 0000000000..8bf9587b94 --- /dev/null +++ b/projects/hip/tests/hipify-clang/cuComplex/cuComplex_Julia.cu @@ -0,0 +1,58 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// CHECK: #include +// 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