From 6f501f60948656a7c1c3e58815d554d374227d19 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 24 Apr 2018 22:04:04 +0300 Subject: [PATCH 1/5] [HIPIFY] Sync with HIP (Execution Control, Surfaces, Memory) Execution Control [deprecated since 7.0]: cudaConfigureCall -> hipConfigureCall cudaLaunch -> hipLaunchByPtr cudaSetupArgument -> hipSetupArgument Surfaces (partially): cudaSurfaceObject_t -> hipSurfaceObject_t cudaSurfaceBoundaryMode -> hipSurfaceBoundaryMode cudaCreateSurfaceObject -> hipCreateSurfaceObject cudaDestroySurfaceObject -> hipDestroySurfaceObject Memory Management: cudaMemset2DAsync -> hipMemset2DAsync --- ..._Runtime_API_functions_supported_by_HIP.md | 22 +++++++++---------- hipify-clang/src/CUDA2HipMap.cpp | 21 +++++++++--------- 2 files changed, 22 insertions(+), 21 deletions(-) diff --git a/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index 84fd0f469d..65528da7fb 100644 --- a/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md @@ -102,9 +102,9 @@ | **CUDA** | **HIP** | |-----------------------------------------------------------|-------------------------------| -| `cudaConfigureCall` | | -| `cudaLaunch` | | -| `cudaSetupArgument` | | +| `cudaConfigureCall` | `hipConfigureCall` | +| `cudaLaunch` | `hipLaunchByPtr` | +| `cudaSetupArgument` | `hipSetupArgument` | ## **9. Memory Management** @@ -158,7 +158,7 @@ | `cudaMemcpyToSymbolAsync` | `hipMemcpyToSymbolAsync` | | `cudaMemset` | `hipMemset` | | `cudaMemset2D` | `hipMemset2D` | -| `cudaMemset2DAsync` | | +| `cudaMemset2DAsync` | `hipMemset2DAsync` | | `cudaMemset3D` | | | `cudaMemset3DAsync` | | | `cudaMemsetAsync` | `hipMemsetAsync` | @@ -338,8 +338,8 @@ | **CUDA** | **HIP** | |-----------------------------------------------------------|-------------------------------| -| `cudaCreateSurfaceObject` | | -| `cudaDestroySurfaceObject` | | +| `cudaCreateSurfaceObject` | `hipCreateSurfaceObject` | +| `cudaDestroySurfaceObject` | `hipDestroySurfaceObject` | | `cudaGetSurfaceObjectResourceDesc` | | ## **27. Version Management** @@ -675,10 +675,10 @@ | 0 |*`cudaSharedMemBankSizeDefault`* |*`hipSharedMemBankSizeDefault`* | | 1 |*`cudaSharedMemBankSizeFourByte`* |*`hipSharedMemBankSizeFourByte`* | | 2 |*`cudaSharedMemBankSizeEightByte`* |*`hipSharedMemBankSizeEightByte`* | -| enum |***`cudaSurfaceBoundaryMode`*** | | -| 0 |*`cudaBoundaryModeZero`* | | -| 1 |*`cudaBoundaryModeClamp`* | | -| 2 |*`cudaBoundaryModeTrap`* | | +| enum |***`cudaSurfaceBoundaryMode`*** |***`hipSurfaceBoundaryMode`*** | +| 0 |*`cudaBoundaryModeZero`* |*`hipBoundaryModeZero`* | +| 1 |*`cudaBoundaryModeClamp`* |*`hipBoundaryModeClamp`* | +| 2 |*`cudaBoundaryModeTrap`* |*`hipBoundaryModeTrap`* | | enum |***`cudaSurfaceFormatMode`*** | | | 0 |*`cudaFormatModeForced`* | | | 1 |*`cudaFormatModeAuto`* | | @@ -742,7 +742,7 @@ | typedef | `cudaOutputMode_t` | | | typedef | `cudaStream_t` | `hipStream_t` | | typedef | `cudaStreamCallback_t` | `hipStreamCallback_t` | -| typedef | `cudaSurfaceObject_t` | | +| typedef | `cudaSurfaceObject_t` | `hipSurfaceObject_t` | | typedef | `cudaTextureObject_t` | `hipTextureObject_t` | | typedef | `CUuuid_stcudaUUID_t` | | | define | `CUDA_IPC_HANDLE_SIZE` | | diff --git a/hipify-clang/src/CUDA2HipMap.cpp b/hipify-clang/src/CUDA2HipMap.cpp index 47358802e9..ddde96e846 100644 --- a/hipify-clang/src/CUDA2HipMap.cpp +++ b/hipify-clang/src/CUDA2HipMap.cpp @@ -261,12 +261,13 @@ const std::map CUDA_TYPE_NAME_MAP{ // typedefs {"cudaTextureObject_t", {"hipTextureObject_t", CONV_TEX, API_RUNTIME}}, + {"cudaSurfaceObject_t", {"hipSurfaceObject_t", CONV_SURFACE, API_RUNTIME}}, // enums {"cudaResourceType", {"hipResourceType", CONV_TEX, API_RUNTIME}}, // API_Driver ANALOGUE (CUresourcetype) {"cudaResourceViewFormat", {"hipResourceViewFormat", CONV_TEX, API_RUNTIME}}, // API_Driver ANALOGUE (CUresourceViewFormat) {"cudaTextureAddressMode", {"hipTextureAddressMode", CONV_TEX, API_RUNTIME}}, - {"cudaSurfaceBoundaryMode", {"hipSurfaceBoundaryMode", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}}, + {"cudaSurfaceBoundaryMode", {"hipSurfaceBoundaryMode", CONV_SURFACE, API_RUNTIME}}, {"cudaSurfaceFormatMode", {"hipSurfaceFormatMode", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}}, @@ -1484,7 +1485,7 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cudaMemset", {"hipMemset", CONV_MEM, API_RUNTIME}}, {"cudaMemsetAsync", {"hipMemsetAsync", CONV_MEM, API_RUNTIME}}, {"cudaMemset2D", {"hipMemset2D", CONV_MEM, API_RUNTIME}}, - {"cudaMemset2DAsync", {"hipMemset2DAsync", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}}, + {"cudaMemset2DAsync", {"hipMemset2DAsync", CONV_MEM, API_RUNTIME}}, {"cudaMemset3D", {"hipMemset3D", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}}, {"cudaMemset3DAsync", {"hipMemset3DAsync", CONV_MEM, API_RUNTIME, HIP_UNSUPPORTED}}, @@ -1745,9 +1746,9 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cudaSetDoubleForHost", {"hipSetDoubleForHost", CONV_EXEC, API_RUNTIME, HIP_UNSUPPORTED}}, // Execution Control [deprecated since 7.0] - {"cudaConfigureCall", {"hipConfigureCall", CONV_EXEC, API_RUNTIME, HIP_UNSUPPORTED}}, - {"cudaLaunch", {"hipLaunch", CONV_EXEC, API_RUNTIME, HIP_UNSUPPORTED}}, - {"cudaSetupArgument", {"hipSetupArgument", CONV_EXEC, API_RUNTIME, HIP_UNSUPPORTED}}, + {"cudaConfigureCall", {"hipConfigureCall", CONV_EXEC, API_RUNTIME}}, + {"cudaLaunch", {"hipLaunchByPtr", CONV_EXEC, API_RUNTIME}}, + {"cudaSetupArgument", {"hipSetupArgument", CONV_EXEC, API_RUNTIME}}, // Version Management {"cudaDriverGetVersion", {"hipDriverGetVersion", CONV_VERSION, API_RUNTIME}}, @@ -1889,17 +1890,17 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cudaGetSurfaceReference", {"hipGetSurfaceReference", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}}, // enum cudaSurfaceBoundaryMode - {"cudaBoundaryModeZero", {"hipBoundaryModeZero", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"cudaBoundaryModeClamp", {"hipBoundaryModeClamp", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"cudaBoundaryModeTrap", {"hipBoundaryModeTrap", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}}, + {"cudaBoundaryModeZero", {"hipBoundaryModeZero", CONV_SURFACE, API_RUNTIME}}, + {"cudaBoundaryModeClamp", {"hipBoundaryModeClamp", CONV_SURFACE, API_RUNTIME}}, + {"cudaBoundaryModeTrap", {"hipBoundaryModeTrap", CONV_SURFACE, API_RUNTIME}}, // enum cudaSurfaceFormatMode {"cudaFormatModeForced", {"hipFormatModeForced", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}}, {"cudaFormatModeAuto", {"hipFormatModeAuto", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}}, // Surface Object Management - {"cudaCreateSurfaceObject", {"hipCreateSurfaceObject", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}}, - {"cudaDestroySurfaceObject", {"hipDestroySurfaceObject", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}}, + {"cudaCreateSurfaceObject", {"hipCreateSurfaceObject", CONV_SURFACE, API_RUNTIME}}, + {"cudaDestroySurfaceObject", {"hipDestroySurfaceObject", CONV_SURFACE, API_RUNTIME}}, {"cudaGetSurfaceObjectResourceDesc", {"hipGetSurfaceObjectResourceDesc", CONV_SURFACE, API_RUNTIME, HIP_UNSUPPORTED}}, // Inter-Process Communications (IPC) From 6411ca1f6d0e24d41ef182614da71e378bd14dff Mon Sep 17 00:00:00 2001 From: Lakhan Singh Date: Thu, 3 May 2018 09:27:50 +0530 Subject: [PATCH 2/5] Null checks added for hipmallocpitch and hipmemcpy apis --- src/hip_memory.cpp | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index 3219dac3db..48a8da81ca 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -339,6 +339,11 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) { // width in bytes hipError_t ihipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height, size_t depth) { hipError_t hip_status = hipSuccess; + if(ptr==NULL) + { + hip_status=hipErrorInvalidValue; + return hip_status; + } // hardcoded 128 bytes *pitch = ((((int)width - 1) / 128) + 1) * 128; const size_t sizeBytes = (*pitch) * height; @@ -1132,7 +1137,11 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind hc::completion_future marker; hipError_t e = hipSuccess; - + if(dst==NULL || src==NULL) + { + e=hipErrorInvalidValue; + return e; + } try { stream->locked_copySync(dst, src, sizeBytes, kind); } catch (ihipException& ex) { From fe421c89b23c275376c3ec107ed30f2b3188ccb3 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 3 May 2018 11:33:40 +0300 Subject: [PATCH 3/5] [HIPIFY] Initial cuDNN support - based on https://github.com/ROCmSoftwarePlatform/hipDNN. - lit testing was supplemented with CUDA_DNN_ROOT_DIR for cuDNN testing. - single cuDNN test was added. --- hipify-clang/src/CUDA2HipMap.cpp | 80 ++++++++++++++++++++++++++++++ hipify-clang/src/HipifyAction.cpp | 4 ++ hipify-clang/src/HipifyAction.h | 1 + hipify-clang/src/Statistics.cpp | 2 +- hipify-clang/src/Statistics.h | 1 + tests/hipify-clang/lit.cfg | 4 +- tests/hipify-clang/lit.site.cfg.in | 1 + 7 files changed, 91 insertions(+), 2 deletions(-) diff --git a/hipify-clang/src/CUDA2HipMap.cpp b/hipify-clang/src/CUDA2HipMap.cpp index ddde96e846..062018be3a 100644 --- a/hipify-clang/src/CUDA2HipMap.cpp +++ b/hipify-clang/src/CUDA2HipMap.cpp @@ -404,6 +404,9 @@ const std::map CUDA_INCLUDE_MAP{ {"curand_precalc.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, {"curand_uniform.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + // CUDNN includes + {"cudnn.h", {"hipDNN.h", CONV_INCLUDE_CUDA_MAIN_H, API_RAND}}, + // HIP includes // TODO: uncomment this when hip/cudacommon.h will be renamed to hip/hipcommon.h // {"cudacommon.h", {"hipcommon.h", CONV_INCLUDE, API_RUNTIME}}, @@ -2881,6 +2884,83 @@ const std::map CUDA_IDENTIFIER_MAP{ {"curand_poisson4", {"hiprand_poisson4", CONV_DEVICE_FUNC, API_RAND}}, {"curand_Philox4x32_10", {"hiprand_Philox4x32_10", CONV_DEVICE_FUNC, API_RAND, HIP_UNSUPPORTED}}, // unchanged function names: skipahead, skipahead_sequence, skipahead_subsequence + + ///////////////////////////// cuDNN ///////////////////////////// + {"cudnnContext", {"hipdnnContext", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnHandle_t", {"hipdnnHandle_t", CONV_TYPE, API_DNN}}, + {"cudnnStatus_t", {"hipdnnStatus_t", CONV_TYPE, API_DNN}}, + {"CUDNN_STATUS_SUCCESS", {"HIPDNN_STATUS_SUCCESS", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 + {"CUDNN_STATUS_NOT_INITIALIZED", {"HIPDNN_STATUS_NOT_INITIALIZED", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 + {"CUDNN_STATUS_ALLOC_FAILED", {"HIPDNN_STATUS_ALLOC_FAILED", CONV_NUMERIC_LITERAL, API_DNN}}, // 2 + {"CUDNN_STATUS_BAD_PARAM", {"HIPDNN_STATUS_BAD_PARAM", CONV_NUMERIC_LITERAL, API_DNN}}, // 3 + {"CUDNN_STATUS_INTERNAL_ERROR", {"HIPDNN_STATUS_INTERNAL_ERROR", CONV_NUMERIC_LITERAL, API_DNN}}, // 4 + {"CUDNN_STATUS_INVALID_VALUE", {"HIPDNN_STATUS_INVALID_VALUE", CONV_NUMERIC_LITERAL, API_DNN}}, // 5 + {"CUDNN_STATUS_ARCH_MISMATCH", {"HIPDNN_STATUS_ARCH_MISMATCH", CONV_NUMERIC_LITERAL, API_DNN}}, // 6 + {"CUDNN_STATUS_MAPPING_ERROR", {"HIPDNN_STATUS_MAPPING_ERROR", CONV_NUMERIC_LITERAL, API_DNN}}, // 7 + {"CUDNN_STATUS_EXECUTION_FAILED", {"HIPDNN_STATUS_EXECUTION_FAILED", CONV_NUMERIC_LITERAL, API_DNN}}, // 8 + {"CUDNN_STATUS_NOT_SUPPORTED", {"HIPDNN_STATUS_NOT_SUPPORTED", CONV_NUMERIC_LITERAL, API_DNN}}, // 9 + {"CUDNN_STATUS_LICENSE_ERROR", {"HIPDNN_STATUS_LICENSE_ERROR", CONV_NUMERIC_LITERAL, API_DNN}}, // 10 + {"CUDNN_STATUS_RUNTIME_PREREQUISITE_MISSING", {"HIPDNN_STATUS_RUNTIME_PREREQUISITE_MISSING", CONV_NUMERIC_LITERAL, API_DNN}}, // 11 + {"CUDNN_STATUS_RUNTIME_IN_PROGRESS", {"HIPDNN_STATUS_RUNTIME_IN_PROGRESS", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 12 + {"CUDNN_STATUS_RUNTIME_FP_OVERFLOW", {"HIPDNN_STATUS_RUNTIME_FP_OVERFLOW", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 13 + {"cudnnRuntimeTag_t", {"hipdnnRuntimeTag_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnTensorDescriptor_t", {"hipdnnTensorDescriptor_t", CONV_TYPE, API_DNN}}, + {"cudnnConvolutionDescriptor_t", {"hipdnnConvolutionDescriptor_t", CONV_TYPE, API_DNN}}, + {"cudnnConvolutionMode_t", {"hipdnnConvolutionMode_t", CONV_TYPE, API_DNN}}, + {"CUDNN_CONVOLUTION", {"HIPDNN_CONVOLUTION", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 + {"CUDNN_CROSS_CORRELATION", {"HIPDNN_CROSS_CORRELATION", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 + {"cudnnTensorFormat_t", {"hipdnnTensorFormat_t", CONV_TYPE, API_DNN}}, + {"CUDNN_TENSOR_NCHW", {"HIPDNN_TENSOR_NCHW", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 + {"CUDNN_TENSOR_NHWC", {"HIPDNN_TENSOR_NHWC", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 + {"CUDNN_TENSOR_NCHW_VECT_C", {"HIPDNN_TENSOR_NCHW_VECT_C", CONV_NUMERIC_LITERAL, API_DNN}}, // 2 + {"cudnnDataType_t", {"hipdnnDataType_t", CONV_TYPE, API_DNN}}, + {"CUDNN_DATA_FLOAT", {"HIPDNN_DATA_FLOAT", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 + {"CUDNN_DATA_DOUBLE", {"HIPDNN_DATA_DOUBLE", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 + {"CUDNN_DATA_HALF", {"HIPDNN_DATA_HALF", CONV_NUMERIC_LITERAL, API_DNN}}, // 2 + {"CUDNN_DATA_INT8", {"HIPDNN_DATA_INT8", CONV_NUMERIC_LITERAL, API_DNN}}, // 3 + {"CUDNN_DATA_INT32", {"HIPDNN_DATA_INT32", CONV_NUMERIC_LITERAL, API_DNN}}, // 4 + {"CUDNN_DATA_INT8x4", {"HIPDNN_DATA_INT8x4", CONV_NUMERIC_LITERAL, API_DNN}}, // 5 + {"CUDNN_DATA_UINT8", {"HIPDNN_DATA_UINT8", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 6 + {"CUDNN_DATA_UINT8x4", {"HIPDNN_DATA_UINT8x4", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 7 + + {"cudnnConvolutionFwdAlgo_t", {"hipdnnConvolutionFwdAlgo_t", CONV_TYPE, API_DNN}}, + {"CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM", {"HIPDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_GEMM", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 + {"CUDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM", {"HIPDNN_CONVOLUTION_FWD_ALGO_IMPLICIT_PRECOMP_GEMM", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 + {"CUDNN_CONVOLUTION_FWD_ALGO_GEMM", {"HIPDNN_CONVOLUTION_FWD_ALGO_GEMM", CONV_NUMERIC_LITERAL, API_DNN}}, // 2 + {"CUDNN_CONVOLUTION_FWD_ALGO_DIRECT", {"HIPDNN_CONVOLUTION_FWD_ALGO_DIRECT", CONV_NUMERIC_LITERAL, API_DNN}}, // 3 + {"CUDNN_CONVOLUTION_FWD_ALGO_FFT", {"HIPDNN_CONVOLUTION_FWD_ALGO_FFT", CONV_NUMERIC_LITERAL, API_DNN}}, // 4 + {"CUDNN_CONVOLUTION_FWD_ALGO_FFT_TILING", {"HIPDNN_CONVOLUTION_FWD_ALGO_FFT_TILING", CONV_NUMERIC_LITERAL, API_DNN}}, // 5 + {"CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD", {"HIPDNN_CONVOLUTION_FWD_ALGO_WINOGRAD", CONV_NUMERIC_LITERAL, API_DNN}}, // 6 + {"CUDNN_CONVOLUTION_FWD_ALGO_WINOGRAD_NONFUSED", {"HIPDNN_CONVOLUTION_FWD_ALGO_WINOGRAD_NONFUSED", CONV_NUMERIC_LITERAL, API_DNN}}, // 7 + {"CUDNN_CONVOLUTION_FWD_ALGO_COUNT", {"HIPDNN_CONVOLUTION_FWD_ALGO_COUNT", CONV_NUMERIC_LITERAL, API_DNN}}, // 8 + + {"cudnnConvolutionFwdPreference_t", {"hipdnnConvolutionFwdPreference_t", CONV_TYPE, API_DNN}}, + {"CUDNN_CONVOLUTION_FWD_NO_WORKSPACE", {"HIPDNN_CONVOLUTION_FWD_NO_WORKSPACE", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 + {"CUDNN_CONVOLUTION_FWD_PREFER_FASTEST", {"HIPDNN_CONVOLUTION_FWD_PREFER_FASTEST", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 + {"CUDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT", {"HIPDNN_CONVOLUTION_FWD_SPECIFY_WORKSPACE_LIMIT", CONV_NUMERIC_LITERAL, API_DNN}}, // 2 + + {"cudnnFilterDescriptor_t", {"hipdnnFilterDescriptor_t", CONV_TYPE, API_DNN}}, + + {"cudnnGetVersion", {"hipdnnGetVersion", CONV_VERSION, API_DNN}}, + {"cudnnGetCudartVersion", {"hipdnnGetCudartVersion", CONV_VERSION, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetErrorString", {"hipdnnGetErrorString", CONV_ERROR, API_DNN}}, + {"cudnnCreate", {"hipdnnCreate", CONV_MATH_FUNC, API_DNN}}, + {"cudnnCreateTensorDescriptor", {"hipdnnCreateTensorDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnSetTensor4dDescriptor", {"hipdnnSetTensor4dDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnSetConvolution2dDescriptor", {"hipdnnSetConvolution2dDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnGetConvolution2dForwardOutputDim", {"hipdnnGetConvolution2dForwardOutputDim", CONV_MATH_FUNC, API_DNN}}, + {"cudnnCreateFilterDescriptor", {"hipdnnCreateFilterDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnSetFilter4dDescriptor", {"hipdnnSetFilter4dDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnCreateConvolutionDescriptor", {"hipdnnCreateConvolutionDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnGetConvolutionForwardAlgorithm", {"hipdnnGetConvolutionForwardAlgorithm", CONV_MATH_FUNC, API_DNN}}, + {"cudnnConvolutionForward", {"hipdnnConvolutionForward", CONV_MATH_FUNC, API_DNN}}, + {"cudnnGetConvolutionForwardWorkspaceSize", {"hipdnnGetConvolutionForwardWorkspaceSize", CONV_MATH_FUNC, API_DNN}}, + {"cudnnDestroyTensorDescriptor", {"hipdnnDestroyTensorDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnDestroyConvolutionDescriptor", {"hipdnnDestroyConvolutionDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnDestroyFilterDescriptor", {"hipdnnDestroyFilterDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnDestroyFilterDescriptor", {"hipdnnDestroyFilterDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnDestroy", {"hipdnnDestroy", CONV_MATH_FUNC, API_DNN}}, + }; const std::map& CUDA_RENAMES_MAP() { diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp index 76efe9bf85..7e5ff4357d 100644 --- a/hipify-clang/src/HipifyAction.cpp +++ b/hipify-clang/src/HipifyAction.cpp @@ -160,6 +160,10 @@ bool HipifyAction::Exclude(const hipCounter & hipToken) { insertedRANDHeader = true; return false; } + case API_DNN: + if (insertedDNNHeader) { return true; } + insertedDNNHeader = true; + return false; default: return false; } diff --git a/hipify-clang/src/HipifyAction.h b/hipify-clang/src/HipifyAction.h index 8c2195b28d..af75f93063 100644 --- a/hipify-clang/src/HipifyAction.h +++ b/hipify-clang/src/HipifyAction.h @@ -27,6 +27,7 @@ private: bool insertedBLASHeader = false; bool insertedRANDHeader = false; bool insertedRAND_kernelHeader = false; + bool insertedDNNHeader = false; bool firstHeader = false; bool pragmaOnce = false; clang::SourceLocation firstHeaderLoc; diff --git a/hipify-clang/src/Statistics.cpp b/hipify-clang/src/Statistics.cpp index 4c5664a6b0..13d4c3effc 100644 --- a/hipify-clang/src/Statistics.cpp +++ b/hipify-clang/src/Statistics.cpp @@ -14,7 +14,7 @@ const char *counterNames[NUM_CONV_TYPES] = { }; const char *apiNames[NUM_API_TYPES] = { - "CUDA Driver API", "CUDA RT API", "CUBLAS API", "CURAND API" + "CUDA Driver API", "CUDA RT API", "CUBLAS API", "CURAND API", "CUDNN API" }; namespace { diff --git a/hipify-clang/src/Statistics.h b/hipify-clang/src/Statistics.h index 81be7b09a8..77bf9b0518 100644 --- a/hipify-clang/src/Statistics.h +++ b/hipify-clang/src/Statistics.h @@ -56,6 +56,7 @@ enum ApiTypes { API_RUNTIME, API_BLAS, API_RAND, + API_DNN, API_LAST }; constexpr int NUM_API_TYPES = (int) ApiTypes::API_LAST; diff --git a/tests/hipify-clang/lit.cfg b/tests/hipify-clang/lit.cfg index 104cee8311..98961fe166 100644 --- a/tests/hipify-clang/lit.cfg +++ b/tests/hipify-clang/lit.cfg @@ -57,6 +57,8 @@ else: run_test_ext = ".sh" clang_args += " -isystem'%s'/samples/common/inc" -config.substitutions.append(("%cuda_args", clang_args % (config.cuda_root, config.cuda_sdk_root))) +clang_args += " -I'%s'/include" + +config.substitutions.append(("%cuda_args", clang_args % (config.cuda_root, config.cuda_sdk_root, config.cuda_dnn_root))) config.substitutions.append(("hipify", '"' + hipify_path + "/hipify-clang" + '"')) config.substitutions.append(("%run_test", '"' + config.test_source_root + "/run_test" + run_test_ext + '"')) diff --git a/tests/hipify-clang/lit.site.cfg.in b/tests/hipify-clang/lit.site.cfg.in index e52db7b026..99aa92b49f 100644 --- a/tests/hipify-clang/lit.site.cfg.in +++ b/tests/hipify-clang/lit.site.cfg.in @@ -4,6 +4,7 @@ import os config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@" config.obj_root = "@CMAKE_CURRENT_BINARY_DIR@" config.cuda_root = "@CUDA_TOOLKIT_ROOT_DIR@" +config.cuda_dnn_root = "@CUDA_DNN_ROOT_DIR@" if sys.platform in ['win32']: config.cuda_sdk_root = "@CUDA_SDK_ROOT_DIR@" if not config.cuda_sdk_root or config.cuda_sdk_root == "CUDA_SDK_ROOT_DIR-NOTFOUND": From 054c3f71f03ca91d24d8409e6ecd9f259d55d644 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 4 May 2018 10:50:18 +0300 Subject: [PATCH 4/5] [HIPIFY][doc] Readme.md update + supported CUDA version to LLVM version correspondence table is added. + Test section is rewritten. + Windows support is added. --- hipify-clang/README.md | 136 ++++++++++++++++++++++++++++++++++++----- 1 file changed, 121 insertions(+), 15 deletions(-) diff --git a/hipify-clang/README.md b/hipify-clang/README.md index 563440e709..c6f59a0251 100644 --- a/hipify-clang/README.md +++ b/hipify-clang/README.md @@ -17,41 +17,146 @@ ### Dependencies -`hipify-clang` requires clang+llvm of at least version 3.8. +`hipify-clang` requires: +1. LLVM+CLANG of at least version 3.8.0, latest stable release is 6.0.0. +2. CUDA at least version 7.5, latest supported release is 9.0. -In most cases, you can get a suitable version of clang+llvm with your package manager. +| **LLVM release version** | **CUDA latest supported version** | +|:------------------------:|:---------------------------------:| +| 3.8.0 | 7.5 | +| 3.8.1 | 7.5 | +| 3.9.0 | 7.5 | +| 3.9.1 | 7.5 | +| 4.0.0 | 8.0 | +| 4.0.1 | 8.0 | +| 5.0.0 | 8.0 | +| 5.0.1 | 8.0 | +| 6.0.0 | 9.0 | -Failing that, you can [download a release archive](http://releases.llvm.org/), extract it somewhere, and set -[CMAKE_PREFIX_PATH](https://cmake.org/cmake/help/v3.0/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it. +In most cases, you can get a suitable version of LLVM+CLANG with your package manager. + +Failing that or having multiple versions of LLVM, you can [download a release archive](http://releases.llvm.org/), build or install it, and set +[CMAKE_PREFIX_PATH](https://cmake.org/cmake/help/v3.0/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it; for instance: `-DCMAKE_PREFIX_PATH=f:\LLVM\6.0.0\dist` ### Build + Assuming this repository is at `./HIP`: ```shell -mkdir build inst - +cd hipify-clang +mkdir build dist cd build + cmake \ - -DCMAKE_INSTALL_PREFIX=../inst \ + -DCMAKE_INSTALL_PREFIX=../dist \ -DCMAKE_BUILD_TYPE=Release \ - -DBUILD_HIPIFY_CLANG=ON \ - ../HIP + .. make -j install ``` +On Windows the following option should be specified for `cmake` at first place: `-G "Visual Studio 15 2017 Win64"` and after `cmake` the generated `hipify-clang.sln` should be built by `Visual Studio 15 2017` instead of `make`. -The binary can then be found at `./inst/bin/hipify-clang`. +Debug build type `-DCMAKE_BUILD_TYPE=Debug` is also supported and tested, `LLVM+CLANG` should be built in `Debug` mode as well. +64 bit build mode `-Thost=x64` is supported as well, `LLVM+CLANG` should be built (installed) in 64bit mode as well. + +The binary can then be found at `./dist/bin/hipify-clang`. ### Test `hipify-clang` has unit tests using LLVM [`lit`](https://llvm.org/docs/CommandGuide/lit.html)/[`FileCheck`](https://llvm.org/docs/CommandGuide/FileCheck.html). +**LLVM+CLANG should be built from sources, Pre-Built Binaries are not exhaustive for testing.** + To run it: -1. Ensure `lit` and `FileCheck` are installed - these are distributed with LLVM. -2. Ensure `socat` is installed - your distro almost certainly has a package for this. -3. Build with the `HIPIFY_CLANG_TESTS` option turned on. -4. `make test-hipify` +1. Download [`LLVM`](http://releases.llvm.org/6.0.0/llvm-6.0.0.src.tar.xz)+[`CLANG`](http://releases.llvm.org/6.0.0/cfe-6.0.0.src.tar.xz) sources. +2. Build [`LLVM+CLANG`](http://llvm.org/docs/CMake.html). + For instance: + ```shell + cd llvm + mkdir build dist + cd build + + cmake \ + -DCMAKE_INSTALL_PREFIX=../dist \ + -DLLVM_SOURCE_DIR=../llvm \ + -DCMAKE_BUILD_TYPE=Debug \ + -Thost=x64 \ + ../llvm + + make -j install + ``` + On Windows the following option should be specified for `cmake` at first place: `-G "Visual Studio 15 2017 Win64"` and after `cmake` the generated `LLVM.sln` should be built by `Visual Studio 15 2017` instead of `make`. + +3. Ensure [`CUDA`](https://developer.nvidia.com/cuda-toolkit-archive) of minimum version 7.5 is installed. + + * Having multiple CUDA installations in order to choose a concrete version the `DCUDA_TOOLKIT_ROOT_DIR` option should be specified: + + `-DCUDA_TOOLKIT_ROOT_DIR="C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0"` + + * On Windows `CUDA_SDK_ROOT_DIR` option should be specified as well: + + `-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v9.0"` + +4. Ensure [`cuDNN`](https://developer.nvidia.com/rdp/cudnn-archive) of version corresponding to CUDA's version is installed. + + * Path to cuDNN should be specified by the `CUDA_DNN_ROOT_DIR` option: + + `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.1` + +5. Ensure [`python`](https://www.python.org/downloads) of minimum required version 2.7 is installed. +6. Ensure `lit` and `FileCheck` are installed - these are distributed with LLVM. + * installing `lit` into `python` might be required: + + `python f:/LLVM/6.0.0/llvm/utils/lit/setup.py install`, + + where `f:/LLVM/6.0.0/llvm` is LLVM sources root directory. + + * Starting with LLVM 6.0.0 path to llvm-lit.py script should be specified by the `LLVM_EXTERNAL_LIT` option: + + `-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.0/build/Debug/bin/llvm-lit.py`, + + where `f:/LLVM/6.0.0/build/Debug` is LLVM build directory. +7. Build with the `HIPIFY_CLANG_TESTS` option turned on: -DHIPIFY_CLANG_TESTS=1. +8. `make test-hipify` + + On Windows after `cmake` the project `test-hipify` in the generated `hipify-clang.sln` should be built by `Visual Studio 15 2017` instead of `make test-hipify`. + +### Windows + +On Windows the following tested configuration is recommended: + +LLVM 6.0.0 (exact), CUDA 9.0 (exact), cudnn-9.0 (exact), Python 3.6 (min), cmake 3.10 (min), Visual Studio 15.5 2017 (min). + +Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 15 2017`: + +```shell +cmake + -G "Visual Studio 15 2017 Win64" \ + -DHIPIFY_CLANG_TESTS=1 \ + -DCMAKE_BUILD_TYPE=Debug \ + -DCMAKE_INSTALL_PREFIX=../dist \ + -DCMAKE_PREFIX_PATH=f:/LLVM/6.0.0/dist \ + -DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0" \ + -DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v9.0" \ + -DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.1 \ + -DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.0/build/Debug/bin/llvm-lit.py \ + .. +``` +A corresponding successful output: +```shell +-- Found LLVM 6.0.0: +-- - CMake module path: F:/LLVM/6.0.0/dist/lib/cmake/llvm +-- - Include path : F:/LLVM/6.0.0/dist/include +-- - Binary path : F:/LLVM/6.0.0/dist/bin +-- Found PythonInterp: C:/Program Files/Python36/python.exe (found suitable version "3.6.4", minimum required is "2.7") +-- Found lit: C:/Program Files/Python36/Scripts/lit.exe +-- Found FileCheck: F:/LLVM/6.0.0/dist/bin/FileCheck.exe +-- Found CUDA: C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0 (found version "9.0") +-- Configuring done +-- Generating done +-- Build files have been written to: f:/HIP/hipify-clang/build +``` ## Running and using hipify-clang @@ -65,6 +170,7 @@ hipify-clang square.cu -- \ --cuda-path=/opt/cuda \ --cuda-gpu-arch=sm_30 \ -isystem /opt/cuda/samples/common/inc + -I /opt/cuda/cuDNN ``` `hipify-clang` arguments are given first, followed by a separator, and then the arguments you'd pass to `clang` if you @@ -79,5 +185,5 @@ The information contained herein is for informational purposes only, and is subj AMD, the AMD Arrow logo, and combinations thereof are trademarks of Advanced Micro Devices, Inc. Other product names used in this publication are for identification purposes only and may be trademarks of their respective companies. -Copyright (c) 2014-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2014-2018 Advanced Micro Devices, Inc. All rights reserved. From a56b480c5ebeaab7142b86c274f960c5c28ce7af Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 4 May 2018 12:37:15 +0300 Subject: [PATCH 5/5] [HIPIFY][test] add cuDNN test --- .../cuDNN/cudnn_convolution_forward.cu | 254 ++++++++++++++++++ 1 file changed, 254 insertions(+) create mode 100644 tests/hipify-clang/cuDNN/cudnn_convolution_forward.cu diff --git a/tests/hipify-clang/cuDNN/cudnn_convolution_forward.cu b/tests/hipify-clang/cuDNN/cudnn_convolution_forward.cu new file mode 100644 index 0000000000..e58116e22a --- /dev/null +++ b/tests/hipify-clang/cuDNN/cudnn_convolution_forward.cu @@ -0,0 +1,254 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +#include +#include +#include +#include + +// CHECK: #include +#include +// CHECK: #include "hipDNN.h" +#include "cudnn.h" + +// CHECK: hipError_t err = (f); \ +// CHECK: if (err != hipSuccess) { \ + +#define CUDA_CALL(f) { \ + cudaError_t err = (f); \ + if (err != cudaSuccess) { \ + std::cout \ + << " Error occurred: " << err << std::endl; \ + std::exit(1); \ + } \ +} +// CHECK: hipdnnStatus_t err = (f); \ +// CHECK: if (err != HIPDNN_STATUS_SUCCESS) { \ + +#define CUDNN_CALL(f) { \ + cudnnStatus_t err = (f); \ + if (err != CUDNN_STATUS_SUCCESS) { \ + std::cout \ + << " Error occurred: " << err << std::endl; \ + std::exit(1); \ + } \ +} + +__global__ void dev_const(float *px, float k) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + px[tid] = k; +} + +__global__ void dev_iota(float *px) { + int tid = threadIdx.x + blockIdx.x * blockDim.x; + px[tid] = tid; +} + +void print(const float *data, int n, int c, int h, int w) { + std::vector buffer(1 << 20); + // CHECK: CUDA_CALL(hipMemcpy( + CUDA_CALL(cudaMemcpy( + buffer.data(), data, + n * c * h * w * sizeof(float), + // CHECK: hipMemcpyDeviceToHost)); + cudaMemcpyDeviceToHost)); + int a = 0; + for (int i = 0; i < n; ++i) { + for (int j = 0; j < c; ++j) { + std::cout << "n=" << i << ", c=" << j << ":" << std::endl; + for (int k = 0; k < h; ++k) { + for (int l = 0; l < w; ++l) { + std::cout << std::setw(4) << std::right << buffer[a]; + ++a; + } + std::cout << std::endl; + } + } + } + std::cout << std::endl; +} + +int main() { + // CHECK: hipdnnHandle_t cudnn; + cudnnHandle_t cudnn; + // CHECK: CUDNN_CALL(hipdnnCreate(&cudnn)); + CUDNN_CALL(cudnnCreate(&cudnn)); + + // input + const int in_n = 1; + const int in_c = 1; + const int in_h = 5; + const int in_w = 5; + std::cout << "in_n: " << in_n << std::endl; + std::cout << "in_c: " << in_c << std::endl; + std::cout << "in_h: " << in_h << std::endl; + std::cout << "in_w: " << in_w << std::endl; + std::cout << std::endl; + // CHECK: hipdnnTensorDescriptor_t in_desc; + cudnnTensorDescriptor_t in_desc; + // CHECK: CUDNN_CALL(hipdnnCreateTensorDescriptor(&in_desc)); + CUDNN_CALL(cudnnCreateTensorDescriptor(&in_desc)); + // CHECK: CUDNN_CALL(hipdnnSetTensor4dDescriptor( + CUDNN_CALL(cudnnSetTensor4dDescriptor( + // CHECK: in_desc, HIPDNN_TENSOR_NCHW, HIPDNN_DATA_FLOAT, + in_desc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, + in_n, in_c, in_h, in_w)); + + float *in_data; + // CHECK: CUDA_CALL(hipMalloc( + CUDA_CALL(cudaMalloc( + &in_data, in_n * in_c * in_h * in_w * sizeof(float))); + + // filter + const int filt_k = 1; + const int filt_c = 1; + const int filt_h = 2; + const int filt_w = 2; + std::cout << "filt_k: " << filt_k << std::endl; + std::cout << "filt_c: " << filt_c << std::endl; + std::cout << "filt_h: " << filt_h << std::endl; + std::cout << "filt_w: " << filt_w << std::endl; + std::cout << std::endl; + + // CHECK: hipdnnFilterDescriptor_t filt_desc; + cudnnFilterDescriptor_t filt_desc; + // CHECK: CUDNN_CALL(hipdnnCreateFilterDescriptor(&filt_desc)); + CUDNN_CALL(cudnnCreateFilterDescriptor(&filt_desc)); + // CHECK: CUDNN_CALL(hipdnnSetFilter4dDescriptor( + CUDNN_CALL(cudnnSetFilter4dDescriptor( + // CHECK: filt_desc, HIPDNN_DATA_FLOAT, HIPDNN_TENSOR_NCHW, + filt_desc, CUDNN_DATA_FLOAT, CUDNN_TENSOR_NCHW, + filt_k, filt_c, filt_h, filt_w)); + + float *filt_data; + // CUDA_CALL(hipMalloc( + CUDA_CALL(cudaMalloc( + &filt_data, filt_k * filt_c * filt_h * filt_w * sizeof(float))); + + // convolution + const int pad_h = 1; + const int pad_w = 1; + const int str_h = 1; + const int str_w = 1; + const int dil_h = 1; + const int dil_w = 1; + std::cout << "pad_h: " << pad_h << std::endl; + std::cout << "pad_w: " << pad_w << std::endl; + std::cout << "str_h: " << str_h << std::endl; + std::cout << "str_w: " << str_w << std::endl; + std::cout << "dil_h: " << dil_h << std::endl; + std::cout << "dil_w: " << dil_w << std::endl; + std::cout << std::endl; + + // CHECK: hipdnnConvolutionDescriptor_t conv_desc; + cudnnConvolutionDescriptor_t conv_desc; + // CUDNN_CALL(hipdnnCreateConvolutionDescriptor(&conv_desc)); + CUDNN_CALL(cudnnCreateConvolutionDescriptor(&conv_desc)); + // CHECK: CUDNN_CALL(hipdnnSetConvolution2dDescriptor( + CUDNN_CALL(cudnnSetConvolution2dDescriptor( + conv_desc, + pad_h, pad_w, str_h, str_w, dil_h, dil_w, + // CHECK: HIPDNN_CONVOLUTION, HIPDNN_DATA_FLOAT)); + CUDNN_CONVOLUTION, CUDNN_DATA_FLOAT)); + + // output + int out_n; + int out_c; + int out_h; + int out_w; + + // CHECK: CUDNN_CALL(hipdnnGetConvolution2dForwardOutputDim( + CUDNN_CALL(cudnnGetConvolution2dForwardOutputDim( + conv_desc, in_desc, filt_desc, + &out_n, &out_c, &out_h, &out_w)); + + std::cout << "out_n: " << out_n << std::endl; + std::cout << "out_c: " << out_c << std::endl; + std::cout << "out_h: " << out_h << std::endl; + std::cout << "out_w: " << out_w << std::endl; + std::cout << std::endl; + // CHECK: hipdnnTensorDescriptor_t out_desc; + cudnnTensorDescriptor_t out_desc; + // CHECK: CUDNN_CALL(hipdnnCreateTensorDescriptor(&out_desc)); + CUDNN_CALL(cudnnCreateTensorDescriptor(&out_desc)); + // CHECK: CUDNN_CALL(hipdnnSetTensor4dDescriptor( + CUDNN_CALL(cudnnSetTensor4dDescriptor( + // CHECK: out_desc, HIPDNN_TENSOR_NCHW, HIPDNN_DATA_FLOAT, + out_desc, CUDNN_TENSOR_NCHW, CUDNN_DATA_FLOAT, + out_n, out_c, out_h, out_w)); + + float *out_data; + // CHECK: CUDA_CALL(hipMalloc( + CUDA_CALL(cudaMalloc( + &out_data, out_n * out_c * out_h * out_w * sizeof(float))); + + // algorithm + // CHECK: hipdnnConvolutionFwdAlgo_t algo; + cudnnConvolutionFwdAlgo_t algo; + // CHECK: CUDNN_CALL(hipdnnGetConvolutionForwardAlgorithm( + CUDNN_CALL(cudnnGetConvolutionForwardAlgorithm( + cudnn, + in_desc, filt_desc, conv_desc, out_desc, + // CHECK: HIPDNN_CONVOLUTION_FWD_PREFER_FASTEST, 0, &algo)); + CUDNN_CONVOLUTION_FWD_PREFER_FASTEST, 0, &algo)); + + std::cout << "Convolution algorithm: " << algo << std::endl; + std::cout << std::endl; + + // workspace + size_t ws_size; + // CHECK: CUDNN_CALL(hipdnnGetConvolutionForwardWorkspaceSize( + CUDNN_CALL(cudnnGetConvolutionForwardWorkspaceSize( + cudnn, in_desc, filt_desc, conv_desc, out_desc, algo, &ws_size)); + + float *ws_data; + // CHECK: CUDA_CALL(hipMalloc(&ws_data, ws_size)); + CUDA_CALL(cudaMalloc(&ws_data, ws_size)); + + std::cout << "Workspace size: " << ws_size << std::endl; + std::cout << std::endl; + + // perform + float alpha = 1.f; + float beta = 0.f; + // CHECK: hipLaunchKernelGGL(dev_iota, dim3(in_w * in_h), dim3(in_n * in_c), 0, 0, in_data); + // CHECK: hipLaunchKernelGGL(dev_const, dim3(filt_w * filt_h), dim3(filt_k * filt_c), 0, 0, filt_data, 1.f); + dev_iota<<>>(in_data); + dev_const<<>>(filt_data, 1.f); + // CHECK: CUDNN_CALL(hipdnnConvolutionForward( + CUDNN_CALL(cudnnConvolutionForward( + cudnn, + &alpha, in_desc, in_data, filt_desc, filt_data, + conv_desc, algo, ws_data, ws_size, + &beta, out_desc, out_data)); + + // results + std::cout << "in_data:" << std::endl; + print(in_data, in_n, in_c, in_h, in_w); + + std::cout << "filt_data:" << std::endl; + print(filt_data, filt_k, filt_c, filt_h, filt_w); + + std::cout << "out_data:" << std::endl; + print(out_data, out_n, out_c, out_h, out_w); + + // finalizing + // CHECK: CUDA_CALL(hipFree(ws_data)); + CUDA_CALL(cudaFree(ws_data)); + // CHECK: CUDA_CALL(hipFree(out_data)); + CUDA_CALL(cudaFree(out_data)); + // CHECK: CUDNN_CALL(hipdnnDestroyTensorDescriptor(out_desc)); + CUDNN_CALL(cudnnDestroyTensorDescriptor(out_desc)); + // CHECK: CUDNN_CALL(hipdnnDestroyConvolutionDescriptor(conv_desc)); + CUDNN_CALL(cudnnDestroyConvolutionDescriptor(conv_desc)); + // CHECK: CUDA_CALL(hipFree(filt_data)); + CUDA_CALL(cudaFree(filt_data)); + // CHECK: CUDNN_CALL(hipdnnDestroyFilterDescriptor(filt_desc)); + CUDNN_CALL(cudnnDestroyFilterDescriptor(filt_desc)); + // CHECK: CUDA_CALL(hipFree(in_data)); + CUDA_CALL(cudaFree(in_data)); + // CHECK: CUDNN_CALL(hipdnnDestroyTensorDescriptor(in_desc)); + CUDNN_CALL(cudnnDestroyTensorDescriptor(in_desc)); + // CHECK: CUDNN_CALL(hipdnnDestroy(cudnn)); + CUDNN_CALL(cudnnDestroy(cudnn)); + return 0; +}