From 56734306ed3c2fe2c83c2d6440d81be764258251 Mon Sep 17 00:00:00 2001 From: Paul Date: Tue, 24 Apr 2018 14:04:24 -0500 Subject: [PATCH 01/10] Add host and device targets --- hipamd/CMakeLists.txt | 7 +++++-- hipamd/hip-config.cmake.in | 2 +- 2 files changed, 6 insertions(+), 3 deletions(-) diff --git a/hipamd/CMakeLists.txt b/hipamd/CMakeLists.txt index e9a3451222..62676eb7cd 100644 --- a/hipamd/CMakeLists.txt +++ b/hipamd/CMakeLists.txt @@ -215,7 +215,10 @@ if(HIP_PLATFORM STREQUAL "hcc") foreach(TARGET hip_hcc hip_hcc_static hip_device) target_include_directories(${TARGET} SYSTEM INTERFACE $/include>;${HSA_PATH}/include) endforeach() - target_link_libraries(hip_hcc INTERFACE hcc::hccrt;hcc::hc_am) + add_library(host INTERFACE) + target_link_libraries(host INTERFACE hip_hcc) + add_library(device INTERFACE) + target_link_libraries(device INTERFACE host hip_device hcc::hccrt hcc::hc_am) # Generate .hipInfo file(WRITE "${PROJECT_BINARY_DIR}/.hipInfo" ${_buildInfo}) @@ -265,7 +268,7 @@ set(BIN_INSTALL_DIR ${CMAKE_INSTALL_PREFIX}/bin) set(CONFIG_PACKAGE_INSTALL_DIR ${LIB_INSTALL_DIR}/cmake/hip) if(HIP_PLATFORM STREQUAL "hcc") - install(TARGETS hip_hcc_static hip_hcc hip_device EXPORT hip-targets DESTINATION ${LIB_INSTALL_DIR}) + install(TARGETS hip_hcc_static hip_hcc hip_device host device EXPORT hip-targets DESTINATION ${LIB_INSTALL_DIR}) install(EXPORT hip-targets DESTINATION ${CONFIG_PACKAGE_INSTALL_DIR} NAMESPACE hip::) include(CMakePackageConfigHelpers) diff --git a/hipamd/hip-config.cmake.in b/hipamd/hip-config.cmake.in index 7e4468b94a..efcdf708bb 100644 --- a/hipamd/hip-config.cmake.in +++ b/hipamd/hip-config.cmake.in @@ -51,7 +51,7 @@ set_and_check(hip_HIPCONFIG_EXECUTABLE "${hip_BIN_INSTALL_DIR}/hipconfig") find_dependency(hcc) include( "${CMAKE_CURRENT_LIST_DIR}/hip-targets.cmake" ) -set( hip_LIBRARIES hip::hip_hcc) +set( hip_LIBRARIES hip::host hip::device) set( hip_LIBRARY ${hip_LIBRARIES}) set(HIP_INCLUDE_DIR ${hip_INCLUDE_DIR}) From bf9529aaa86bc4660e10c90c58311b00a9ccf2c6 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Fri, 11 May 2018 03:35:10 +0100 Subject: [PATCH 02/10] Add support for the hipFuncGetAttributes interface. --- .../include/hip/hcc_detail/hip_runtime_api.h | 24 ++++++ .../include/hip/hcc_detail/program_state.hpp | 46 +++++++++-- hipamd/src/functional_grid_launch.inl | 4 +- hipamd/src/hip_module.cpp | 79 ++++++++++++++++--- hipamd/src/program_state.cpp | 3 +- .../module/hipFuncGetAttributes.cpp | 53 +++++++++++++ 6 files changed, 189 insertions(+), 20 deletions(-) create mode 100644 hipamd/tests/src/runtimeApi/module/hipFuncGetAttributes.cpp diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 8becab3c9c..fa39a766bc 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -94,6 +94,19 @@ typedef struct ihipModule_t* hipModule_t; typedef struct ihipModuleSymbol_t* hipFunction_t; +struct hipFuncAttributes { + int binaryVersion; + int cacheModeCA; + size_t constSizeBytes; + size_t localSizeBytes; + int maxDynamicSharedSizeBytes; + int maxThreadsPerBlock; + int numRegs; + int preferredShmemCarveout; + int ptxVersion; + size_t sharedSizeBytes; +}; + typedef struct ihipEvent_t* hipEvent_t; enum hipLimit_t { @@ -2222,6 +2235,17 @@ hipError_t hipModuleUnload(hipModule_t module); */ hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, const char* kname); +/** + * @bried Find out attributes for a given function. + * + * @param [out] attr + * @param [in] func + * + * @returns hipSuccess, hipErrorInvalidDeviceFunction + */ + +hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func); + /** * @brief returns device memory pointer and size of the kernel present in the module with symbol @p * name diff --git a/hipamd/include/hip/hcc_detail/program_state.hpp b/hipamd/include/hip/hcc_detail/program_state.hpp index b6ccafb205..ac689fdb89 100644 --- a/hipamd/include/hip/hcc_detail/program_state.hpp +++ b/hipamd/include/hip/hcc_detail/program_state.hpp @@ -22,8 +22,10 @@ THE SOFTWARE. #pragma once +#include #include #include +#include #include #include @@ -46,11 +48,45 @@ struct hash { inline constexpr bool operator==(hsa_agent_t x, hsa_agent_t y) { return x.handle == y.handle; } namespace hip_impl { -struct Kernel_descriptor { - std::uint64_t kernel_object_; - std::uint32_t group_size_; - std::uint32_t private_size_; - std::string name_; +class Kernel_descriptor { + std::uint64_t kernel_object_{}; + amd_kernel_code_t const* kernel_header_{nullptr}; + std::string name_{}; +public: + Kernel_descriptor() = default; + Kernel_descriptor(std::uint64_t kernel_object, const std::string& name) + : kernel_object_{kernel_object}, name_{name} + { + bool supported{false}; + std::uint16_t min_v{UINT16_MAX}; + auto r = hsa_system_major_extension_supported( + HSA_EXTENSION_AMD_LOADER, 1, &min_v, &supported); + + if (r != HSA_STATUS_SUCCESS || !supported) return; + + hsa_ven_amd_loader_1_01_pfn_t tbl{}; + + r = hsa_system_get_major_extension_table( + HSA_EXTENSION_AMD_LOADER, + 1, + sizeof(tbl), + reinterpret_cast(&tbl)); + + if (r != HSA_STATUS_SUCCESS) return; + if (!tbl.hsa_ven_amd_loader_query_host_address) return; + + r = tbl.hsa_ven_amd_loader_query_host_address( + reinterpret_cast(kernel_object_), + reinterpret_cast(&kernel_header_)); + + if (r != HSA_STATUS_SUCCESS) return; + } + Kernel_descriptor(const Kernel_descriptor&) = default; + Kernel_descriptor(Kernel_descriptor&&) = default; + ~Kernel_descriptor() = default; + + Kernel_descriptor& operator=(const Kernel_descriptor&) = default; + Kernel_descriptor& operator=(Kernel_descriptor&&) = default; operator hipFunction_t() const { // TODO: this is awful and only meant for illustration. return reinterpret_cast(const_cast(this)); diff --git a/hipamd/src/functional_grid_launch.inl b/hipamd/src/functional_grid_launch.inl index b555967ebc..9ecad51476 100644 --- a/hipamd/src/functional_grid_launch.inl +++ b/hipamd/src/functional_grid_launch.inl @@ -107,8 +107,8 @@ namespace hip_impl it0->second.cbegin(), it0->second.cend(), [=](const pair& x) { - return x.first.handle == agent.handle; - }); + return x.first == agent; + }); if (it1 == it0->second.cend()) { throw runtime_error{ diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index 84451a17a0..59831958bd 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -27,10 +27,11 @@ THE SOFTWARE. #include "hsa_helpers.hpp" #include "trace_helper.h" +#include #include #include -#include +#include #include #include #include @@ -71,9 +72,8 @@ struct ihipKernArgInfo { map kernelArguments; struct ihipModuleSymbol_t { - uint64_t _object; // The kernel object. - uint32_t _groupSegmentSize; - uint32_t _privateSegmentSize; + uint64_t _object{}; // The kernel object. + amd_kernel_code_t const* _header{}; string _name; // TODO - review for performance cost. Name is just used for debug. }; @@ -179,8 +179,10 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, aql.grid_size_x = globalWorkSizeX; aql.grid_size_y = globalWorkSizeY; aql.grid_size_z = globalWorkSizeZ; - aql.group_segment_size = f->_groupSegmentSize + sharedMemBytes; - aql.private_segment_size = f->_privateSegmentSize; + aql.group_segment_size = + f->_header->workgroup_group_segment_byte_size + sharedMemBytes; + aql.private_segment_size = + f->_header->workitem_private_segment_byte_size; aql.kernel_object = f->_object; aql.setup = 3 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS; aql.header = @@ -444,10 +446,10 @@ hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const ch if (kernel.handle == 0u) return hipErrorNotFound; - (*func)->_object = kernel_object(kernel); - (*func)->_groupSegmentSize = group_size(kernel); - (*func)->_privateSegmentSize = private_size(kernel); - (*func)->_name = name; + // TODO: refactor the whole ihipThisThat, which is a mess and yields the + // below, due to hipFunction_t being a pointer to ihipModuleSymbol_t. + func[0][0] = *static_cast( + Kernel_descriptor{kernel_object(kernel), name}); return hipSuccess; } @@ -471,6 +473,61 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t h return ihipLogStatus(r); } +namespace +{ + inline + hipFuncAttributes make_function_attributes(const amd_kernel_code_t& header) + { + hipFuncAttributes r{}; + + hipDeviceProp_t prop{}; + hipGetDeviceProperties( + &prop, ihipGetTlsDefaultCtx()->getDevice()->_deviceId); + + r.localSizeBytes = header.workitem_private_segment_byte_size; + r.sharedSizeBytes = header.workgroup_group_segment_byte_size; + r.maxDynamicSharedSizeBytes = + prop.sharedMemPerBlock - r.sharedSizeBytes; + r.numRegs = header.workitem_vgpr_count; + r.maxThreadsPerBlock = r.numRegs ? // TODO: proper query. + std::min(prop.maxThreadsPerBlock, 64 * 1024 / r.numRegs) : + prop.maxThreadsPerBlock; + r.binaryVersion = + header.amd_machine_version_major * 10 + + header.amd_machine_version_minor; + r.ptxVersion = prop.major * 10 + prop.minor; // HIP currently presents itself as PTX 3.0. + + return r; + } +} + +hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func) +{ + if (!func) return hipErrorInvalidDeviceFunction; + + const auto it0 = functions().find(reinterpret_cast(func)); + + if (it0 == functions().cend()) return hipErrorInvalidDeviceFunction; + + auto agent = this_agent(); + const auto it1 = find_if( + it0->second.cbegin(), + it0->second.cend(), + [=](const pair& x) { + return x.first == agent; + }); + + if (it1 == it0->second.cend()) return hipErrorInvalidDeviceFunction; + + const auto header = static_cast(it1->second)->_header; + + if (!header) throw runtime_error{"Ill-formed Kernel_descriptor."}; + + *attr = make_function_attributes(*header); + + return hipSuccess; +} + hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) { if (!module) return hipErrorInvalidValue; @@ -487,7 +544,7 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) { (*module)->executable = hip_impl::load_executable( tmp.empty() ? read_elf_file_as_string(image) : tmp, (*module)->executable, this_agent()); - + return (*module)->executable.handle ? hipSuccess : hipErrorUnknown; } diff --git a/hipamd/src/program_state.cpp b/hipamd/src/program_state.cpp index 0c51fe4694..c4478bec2f 100644 --- a/hipamd/src/program_state.cpp +++ b/hipamd/src/program_state.cpp @@ -382,8 +382,7 @@ const unordered_map>>& fu for (auto&& kernel_symbol : it->second) { r[function.first].emplace_back( agent(kernel_symbol), - Kernel_descriptor{kernel_object(kernel_symbol), group_size(kernel_symbol), - private_size(kernel_symbol), it->first}); + Kernel_descriptor{kernel_object(kernel_symbol), it->first}); } } } diff --git a/hipamd/tests/src/runtimeApi/module/hipFuncGetAttributes.cpp b/hipamd/tests/src/runtimeApi/module/hipFuncGetAttributes.cpp new file mode 100644 index 0000000000..7d3eff5f73 --- /dev/null +++ b/hipamd/tests/src/runtimeApi/module/hipFuncGetAttributes.cpp @@ -0,0 +1,53 @@ +/* +Copyright (c) 2015-Present Advanced Micro Devices, Inc. All rights reserved. +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp + * RUN: %t + * HIT_END + */ + +#include +#include +#include "test_common.h" + +__global__ +void fn(float* px, float* py) +{ + bool a[42]; + __shared__ double b[69]; + + for (auto&& x : b) x = *py++; + for (auto&& x : a) x = *px++ > 0.0; + for (auto&& x : a) if (x) *--py = *--px; +} + +int main() { + hipInit(0); + + hipFuncAttributes attr{}; + + auto r = hipFuncGetAttributes(&attr, reinterpret_cast(&fn)); + + if (r != hipSuccess || attr.maxThreadsPerBlock == 0) { + failed("Failed to read attributes."); + } + + passed(); +} From eded014abcb29c814bb0419fa4a134072548e871 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Fri, 11 May 2018 11:31:46 +0100 Subject: [PATCH 03/10] Don't use magic constants, they're evil. Also clarify that the register count cannot be queried at the moment. --- hipamd/src/hip_module.cpp | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index 59831958bd..b174b7ad5b 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -483,14 +483,17 @@ namespace hipDeviceProp_t prop{}; hipGetDeviceProperties( &prop, ihipGetTlsDefaultCtx()->getDevice()->_deviceId); - + // TODO: at the moment there is no way to query the count of registers + // available per CU, therefore we hardcode it to 64 KiRegisters. + prop.regsPerBlock = prop.regsPerBlock ? prop.regsPerBlock : 64 * 1024. + r.localSizeBytes = header.workitem_private_segment_byte_size; r.sharedSizeBytes = header.workgroup_group_segment_byte_size; r.maxDynamicSharedSizeBytes = prop.sharedMemPerBlock - r.sharedSizeBytes; r.numRegs = header.workitem_vgpr_count; - r.maxThreadsPerBlock = r.numRegs ? // TODO: proper query. - std::min(prop.maxThreadsPerBlock, 64 * 1024 / r.numRegs) : + r.maxThreadsPerBlock = r.numRegs ? + std::min(prop.maxThreadsPerBlock, prop.regsPerBlock / r.numRegs) : prop.maxThreadsPerBlock; r.binaryVersion = header.amd_machine_version_major * 10 + From 40a22d235e272261e46cee5e5f2edf254f9d6a58 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Mon, 14 May 2018 17:15:36 +0100 Subject: [PATCH 04/10] Update hip_module.cpp --- hipamd/src/hip_module.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index b174b7ad5b..e4b050af19 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -506,6 +506,7 @@ namespace hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func) { + if (!attr) return hipErrorInvalidValue; if (!func) return hipErrorInvalidDeviceFunction; const auto it0 = functions().find(reinterpret_cast(func)); From 43fca684c80d98180a9aa3f9f3dfc790c3ba111f Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Fri, 18 May 2018 17:50:45 +0100 Subject: [PATCH 05/10] Update hip_module.cpp Typo. --- hipamd/src/hip_module.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index e4b050af19..a6d486b6de 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -485,7 +485,7 @@ namespace &prop, ihipGetTlsDefaultCtx()->getDevice()->_deviceId); // TODO: at the moment there is no way to query the count of registers // available per CU, therefore we hardcode it to 64 KiRegisters. - prop.regsPerBlock = prop.regsPerBlock ? prop.regsPerBlock : 64 * 1024. + prop.regsPerBlock = prop.regsPerBlock ? prop.regsPerBlock : 64 * 1024; r.localSizeBytes = header.workitem_private_segment_byte_size; r.sharedSizeBytes = header.workgroup_group_segment_byte_size; From 513329980428c1bfa600dd428b2c711bab7f6234 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 21 May 2018 11:33:09 +0530 Subject: [PATCH 06/10] Add hipMalloc3D to nvcc detail Change-Id: I8a5654066ed1504e3b05eddbbdebf05fd52aa149 --- hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 92393ab2e0..bd3ffc1cc1 100644 --- a/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -382,6 +382,10 @@ inline static hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, return hipCUDAErrorTohipError(cudaMallocPitch(ptr, pitch, width, height)); } +inline static hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent) { + return hipCUDAErrorTohipError(cudaMalloc3D(pitchedDevPtr, extent)); +} + inline static hipError_t hipFree(void* ptr) { return hipCUDAErrorTohipError(cudaFree(ptr)); } inline static hipError_t hipMallocHost(void** ptr, size_t size) From 92d8e05aa0b2ca4260d3d4147fb3cf14ed571d38 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 21 May 2018 11:35:03 +0530 Subject: [PATCH 07/10] Disable incomplete unit tests that don't work on nvcc path Change-Id: If5823ec96a3b2497a08c46ab802c5a0158271053 --- hipamd/tests/src/deviceLib/hipMathFunctions.cpp | 2 +- hipamd/tests/src/deviceLib/hipTestHalf.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/hipamd/tests/src/deviceLib/hipMathFunctions.cpp b/hipamd/tests/src/deviceLib/hipMathFunctions.cpp index 78e85ba62a..f89bdae149 100644 --- a/hipamd/tests/src/deviceLib/hipMathFunctions.cpp +++ b/hipamd/tests/src/deviceLib/hipMathFunctions.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp + * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc * RUN: %t * HIT_END */ diff --git a/hipamd/tests/src/deviceLib/hipTestHalf.cpp b/hipamd/tests/src/deviceLib/hipTestHalf.cpp index 24a4d6c53e..5a2aac2b29 100644 --- a/hipamd/tests/src/deviceLib/hipTestHalf.cpp +++ b/hipamd/tests/src/deviceLib/hipTestHalf.cpp @@ -18,7 +18,7 @@ THE SOFTWARE. */ /* HIT_START - * BUILD: %t %s ../test_common.cpp + * BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc * RUN: %t * HIT_END */ From 66d05e6fc3e2992302e6c03665701db97d95a272 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 21 May 2018 15:38:44 +0530 Subject: [PATCH 08/10] hipMemcpy returns success if sizeBytes is 0. Fixes SWDEV-153754 & SWDEV-154178. --- hipamd/src/hip_memory.cpp | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index e2202e7860..443e645097 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -1132,15 +1132,19 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) { HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, kind); + hipError_t e = hipSuccess; + + // Return success if number of bytes to copy is 0 + if (sizeBytes == 0) return ihipLogStatus(e); + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); hc::completion_future marker; - hipError_t e = hipSuccess; if(dst==NULL || src==NULL) { e=hipErrorInvalidValue; - return e; + return ihipLogStatus(e); } try { stream->locked_copySync(dst, src, sizeBytes, kind); From 5585072e63aa0fddc3e98e74453b5c32dbdfe9db Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 21 May 2018 18:31:20 +0300 Subject: [PATCH 09/10] [HIPIFY][DNN] support of cuDNN 7.1.3 - finishing --- hipamd/hipify-clang/src/CUDA2HipMap.cpp | 133 +++++++++++++++++++----- 1 file changed, 109 insertions(+), 24 deletions(-) diff --git a/hipamd/hipify-clang/src/CUDA2HipMap.cpp b/hipamd/hipify-clang/src/CUDA2HipMap.cpp index 6a0dfd6ed9..7d10b35e48 100644 --- a/hipamd/hipify-clang/src/CUDA2HipMap.cpp +++ b/hipamd/hipify-clang/src/CUDA2HipMap.cpp @@ -2923,6 +2923,14 @@ const std::map CUDA_IDENTIFIER_MAP{ // defines {"CUDNN_VERSION", {"HIPDNN_VERSION", CONV_NUMERIC_LITERAL, API_DNN}}, // 7000 {"CUDNN_DIM_MAX", {"HIPDNN_DIM_MAX", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 8 + {"CUDNN_LRN_MIN_N", {"HIPDNN_LRN_MIN_N", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1 + {"CUDNN_LRN_MAX_N", {"HIPDNN_LRN_MAX_N", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 16 + {"CUDNN_LRN_MIN_K", {"HIPDNN_LRN_MIN_K", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1e-5 + {"CUDNN_LRN_MIN_BETA", {"HIPDNN_LRN_MIN_BETA", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0.01 + {"CUDNN_BN_MIN_EPSILON", {"HIPDNN_BN_MIN_EPSILON", CONV_NUMERIC_LITERAL, API_DNN}}, // 1e-5 + {"CUDNN_SEV_ERROR_EN", {"HIPDNN_SEV_ERROR_EN", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, + {"CUDNN_SEV_WARNING_EN", {"HIPDNN_SEV_WARNING_EN", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, + {"CUDNN_SEV_INFO_EN", {"HIPDNN_SEV_INFO_EN", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, {"cudnnContext", {"hipdnnContext", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, {"cudnnHandle_t", {"hipdnnHandle_t", CONV_TYPE, API_DNN}}, @@ -2960,11 +2968,15 @@ const std::map CUDA_IDENTIFIER_MAP{ {"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 - {"cudnnErrQueryMode_t", {"hipdnnErrQueryMode_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, {"CUDNN_ERRQUERY_RAWCODE", {"HIPDNN_ERRQUERY_RAWCODE", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0 {"CUDNN_ERRQUERY_NONBLOCKING", {"HIPDNN_ERRQUERY_NONBLOCKING", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1 {"CUDNN_ERRQUERY_BLOCKING", {"HIPDNN_ERRQUERY_BLOCKING", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 2 + {"cudnnSeverity_t", {"hipdnnSeverity_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"CUDNN_SEV_FATAL", {"HIPDNN_SEV_FATAL", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0 + {"CUDNN_SEV_ERROR", {"HIPDNN_SEV_ERROR", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1 + {"CUDNN_SEV_WARNING", {"HIPDNN_SEV_WARNING", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 2 + {"CUDNN_SEV_INFO", {"HIPDNN_SEV_INFO", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 3 {"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 @@ -2976,21 +2988,25 @@ const std::map CUDA_IDENTIFIER_MAP{ {"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 - {"cudnnDeterminism_t", {"hipdnnDeterminism_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, {"CUDNN_NON_DETERMINISTIC", {"HIPDNN_NON_DETERMINISTIC", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0 {"CUDNN_DETERMINISTIC", {"HIPDNN_DETERMINISTIC", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1 + {"cudnnDivNormMode_t", {"hipdnnDivNormMode_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"CUDNN_DIVNORM_PRECOMPUTED_MEANS", {"HIPDNN_DIVNORM_PRECOMPUTED_MEANS", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0 + {"cudnnCTCLossAlgo_t", {"hipdnnCTCLossAlgo_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"CUDNN_CTC_LOSS_ALGO_DETERMINISTIC", {"HIPDNN_CTC_LOSS_ALGO_DETERMINISTIC", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0 + {"CUDNN_CTC_LOSS_ALGO_NON_DETERMINISTIC", {"HIPDNN_CTC_LOSS_ALGO_NON_DETERMINISTIC", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 1 {"cudnnFilterDescriptor_t", {"hipdnnFilterDescriptor_t", CONV_TYPE, API_DNN}}, {"cudnnDropoutDescriptor_t", {"hipdnnDropoutDescriptor_t", CONV_TYPE, API_DNN}}, {"cudnnConvolutionFwdAlgoPerf_t", {"hipdnnConvolutionFwdAlgoPerf_t", CONV_TYPE, API_DNN}}, {"cudnnConvolutionBwdFilterAlgoPerf_t", {"hipdnnConvolutionBwdFilterAlgoPerf_t", CONV_TYPE, API_DNN}}, {"cudnnRNNDescriptor_t", {"hipdnnRNNDescriptor_t", CONV_TYPE, API_DNN}}, + {"cudnnPersistentRNNPlan", {"hipdnnPersistentRNNPlan", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, {"cudnnPersistentRNNPlan_t", {"hipdnnPersistentRNNPlan_t", CONV_TYPE, API_DNN}}, {"cudnnTensorStruct", {"hipdnnTensorStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, {"cudnnConvolutionStruct", {"hipdnnConvolutionStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, @@ -3010,23 +3026,30 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cudnnCTCLossDescriptor_t", {"hipdnnCTCLossDescriptor_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, {"cudnnCTCLossStruct", {"hipdnnCTCLossStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, {"cudnnConvolutionBwdDataAlgoPerf_t", {"hipdnnConvolutionBwdDataAlgoPerf_t", CONV_TYPE, API_DNN}}, + {"cudnnAlgorithmDescriptor_t", {"hipdnnAlgorithmDescriptor_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnAlgorithmStruct", {"hipdnnAlgorithmStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnAlgorithmPerformance_t", {"hipdnnAlgorithmPerformance_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnAlgorithmPerformanceStruct", {"hipdnnAlgorithmPerformanceStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnRNNStruct", {"hipdnnRNNStruct", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnRNNDescriptor_t", {"hipdnnRNNDescriptor_t", CONV_TYPE, API_DNN}}, + {"cudnnAlgorithm_t", {"hipdnnAlgorithm_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnCallback_t", {"hipdnnCallback_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnDebug_t", {"hipdnnDebug_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnLRNMode_t", {"hipdnnLRNMode_t", CONV_TYPE, API_DNN}}, + {"CUDNN_LRN_CROSS_CHANNEL_DIM1", {"HIPDNN_LRN_CROSS_CHANNEL", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 vs 1 {"cudnnRNNInputMode_t", {"hipdnnRNNInputMode_t", CONV_TYPE, API_DNN}}, {"CUDNN_LINEAR_INPUT", {"HIPDNN_LINEAR_INPUT", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_SKIP_INPUT", {"HIPDNN_SKIP_INPUT", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 - {"cudnnDirectionMode_t", {"hipdnnDirectionMode_t", CONV_TYPE, API_DNN}}, {"CUDNN_UNIDIRECTIONAL", {"HIPDNN_UNIDIRECTIONAL", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_BIDIRECTIONAL", {"HIPDNN_BIDIRECTIONAL", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 - {"cudnnMathType_t", {"hipdnnMathType_t", CONV_TYPE, API_DNN}}, {"CUDNN_DEFAULT_MATH", {"HIPDNN_DEFAULT_MATH", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_TENSOR_OP_MATH", {"HIPDNN_TENSOR_OP_MATH", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 - {"cudnnNanPropagation_t", {"hipdnnNanPropagation_t", CONV_TYPE, API_DNN}}, {"CUDNN_NOT_PROPAGATE_NAN", {"HIPDNN_NOT_PROPAGATE_NAN", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_PROPAGATE_NAN", {"HIPDNN_PROPAGATE_NAN", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 - {"cudnnConvolutionBwdDataAlgo_t", {"hipdnnConvolutionBwdDataAlgo_t", CONV_TYPE, API_DNN}}, {"CUDNN_CONVOLUTION_BWD_DATA_ALGO_0", {"HIPDNN_CONVOLUTION_BWD_DATA_ALGO_0", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_CONVOLUTION_BWD_DATA_ALGO_1", {"HIPDNN_CONVOLUTION_BWD_DATA_ALGO_1", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 @@ -3035,7 +3058,6 @@ const std::map CUDA_IDENTIFIER_MAP{ {"CUDNN_CONVOLUTION_BWD_DATA_ALGO_WINOGRAD", {"HIPDNN_CONVOLUTION_BWD_DATA_ALGO_WINOGRAD", CONV_NUMERIC_LITERAL, API_DNN}}, // 4 {"CUDNN_CONVOLUTION_BWD_DATA_ALGO_WINOGRAD_NONFUSED", {"HIPDNN_CONVOLUTION_BWD_DATA_ALGO_WINOGRAD_NONFUSED", CONV_NUMERIC_LITERAL, API_DNN}}, // 5 {"CUDNN_CONVOLUTION_BWD_DATA_ALGO_COUNT", {"HIPDNN_CONVOLUTION_BWD_DATA_ALGO_TRANSPOSE_GEMM", CONV_NUMERIC_LITERAL, API_DNN}}, // 6 - {"cudnnConvolutionBwdFilterAlgo_t", {"hipdnnConvolutionBwdFilterAlgo_t", CONV_TYPE, API_DNN}}, {"CUDNN_CONVOLUTION_BWD_FILTER_ALGO_0", {"HIPDNN_CONVOLUTION_BWD_FILTER_ALGO_0", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_CONVOLUTION_BWD_FILTER_ALGO_1", {"HIPDNN_CONVOLUTION_BWD_FILTER_ALGO_1", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 @@ -3045,24 +3067,20 @@ const std::map CUDA_IDENTIFIER_MAP{ {"CUDNN_CONVOLUTION_BWD_FILTER_ALGO_WINOGRAD_NONFUSED", {"HIPDNN_CONVOLUTION_BWD_FILTER_ALGO_WINOGRAD_NONFUSED", CONV_NUMERIC_LITERAL, API_DNN}}, // 5 {"CUDNN_CONVOLUTION_BWD_FILTER_ALGO_FFT_TILING", {"HIPDNN_CONVOLUTION_BWD_FILTER_ALGO_FFT_TILING", CONV_NUMERIC_LITERAL, API_DNN}}, // 6 {"CUDNN_CONVOLUTION_BWD_FILTER_ALGO_COUNT", {"HIPDNN_CONVOLUTION_BWD_FILTER_ALGO_COUNT", CONV_NUMERIC_LITERAL, API_DNN}}, // 7 - {"cudnnConvolutionBwdFilterPreference_t", {"hipdnnConvolutionBwdFilterPreference_t", CONV_TYPE, API_DNN}}, {"CUDNN_CONVOLUTION_BWD_FILTER_NO_WORKSPACE", {"HIPDNN_CONVOLUTION_BWD_FILTER_NO_WORKSPACE", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_CONVOLUTION_BWD_FILTER_PREFER_FASTEST", {"HIPDNN_CONVOLUTION_BWD_FILTER_PREFER_FASTEST", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 {"CUDNN_CONVOLUTION_BWD_FILTER_SPECIFY_WORKSPACE_LIMIT",{"HIPDNN_CONVOLUTION_BWD_FILTER_SPECIFY_WORKSPACE_LIMIT",CONV_NUMERIC_LITERAL, API_DNN}}, // 2 - {"cudnnRNNAlgo_t", {"hipdnnRNNAlgo_t", CONV_TYPE, API_DNN}}, {"CUDNN_RNN_ALGO_STANDARD", {"HIPDNN_RNN_ALGO_STANDARD", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_RNN_ALGO_PERSIST_STATIC", {"HIPDNN_RNN_ALGO_PERSIST_STATIC", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 {"CUDNN_RNN_ALGO_PERSIST_DYNAMIC", {"HIPDNN_RNN_ALGO_PERSIST_DYNAMIC", CONV_NUMERIC_LITERAL, API_DNN}}, // 2 {"CUDNN_RNN_ALGO_COUNT", {"HIPDNN_RNN_ALGO_COUNT", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 3 - {"cudnnRNNMode_t", {"hipdnnRNNMode_t", CONV_TYPE, API_DNN}}, {"CUDNN_RNN_RELU", {"HIPDNN_RNN_RELU", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_RNN_TANH", {"HIPDNN_RNN_TANH", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 {"CUDNN_LSTM", {"HIPDNN_LSTM", CONV_NUMERIC_LITERAL, API_DNN}}, // 2 {"CUDNN_GRU", {"HIPDNN_GRU", CONV_NUMERIC_LITERAL, API_DNN}}, // 3 - {"cudnnOpTensorOp_t", {"hipdnnOpTensorOp_t", CONV_TYPE, API_DNN}}, {"CUDNN_OP_TENSOR_ADD", {"HIPDNN_OP_TENSOR_ADD", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_OP_TENSOR_MUL", {"HIPDNN_OP_TENSOR_MUL", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 @@ -3070,7 +3088,6 @@ const std::map CUDA_IDENTIFIER_MAP{ {"CUDNN_OP_TENSOR_MAX", {"HIPDNN_OP_TENSOR_MAX", CONV_NUMERIC_LITERAL, API_DNN}}, // 3 {"CUDNN_OP_TENSOR_SQRT", {"HIPDNN_OP_TENSOR_SQRT", CONV_NUMERIC_LITERAL, API_DNN}}, // 4 {"CUDNN_OP_TENSOR_NOT", {"HIPDNN_OP_TENSOR_NOT", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 5 - {"cudnnReduceTensorOp_t", {"hipdnnReduceTensorOp_t", CONV_TYPE, API_DNN}}, {"CUDNN_REDUCE_TENSOR_ADD", {"HIPDNN_REDUCE_TENSOR_ADD", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_REDUCE_TENSOR_MUL", {"HIPDNN_REDUCE_TENSOR_MUL", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 @@ -3081,36 +3098,43 @@ const std::map CUDA_IDENTIFIER_MAP{ {"CUDNN_REDUCE_TENSOR_NORM1", {"HIPDNN_REDUCE_TENSOR_NORM1", CONV_NUMERIC_LITERAL, API_DNN}}, // 6 {"CUDNN_REDUCE_TENSOR_NORM2", {"HIPDNN_REDUCE_TENSOR_NORM2", CONV_NUMERIC_LITERAL, API_DNN}}, // 7 {"CUDNN_REDUCE_TENSOR_MUL_NO_ZEROS", {"HIPDNN_REDUCE_TENSOR_MUL_NO_ZEROS", CONV_NUMERIC_LITERAL, API_DNN}}, // 8 - {"cudnnReduceTensorIndices_t", {"hipdnnReduceTensorIndices_t", CONV_TYPE, API_DNN}}, {"CUDNN_REDUCE_TENSOR_NO_INDICES", {"HIPDNN_REDUCE_TENSOR_NO_INDICES", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_REDUCE_TENSOR_FLATTENED_INDICES", {"HIPDNN_REDUCE_TENSOR_FLATTENED_INDICES", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 - {"cudnnConvolutionBwdDataPreference_t", {"hipdnnConvolutionBwdDataPreference_t", CONV_TYPE, API_DNN}}, {"CUDNN_CONVOLUTION_BWD_DATA_NO_WORKSPACE", {"HIPDNN_CONVOLUTION_BWD_DATA_NO_WORKSPACE", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST", {"HIPDNN_CONVOLUTION_BWD_DATA_PREFER_FASTEST", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 {"CUDNN_CONVOLUTION_BWD_DATA_SPECIFY_WORKSPACE_LIMIT", {"HIPDNN_CONVOLUTION_BWD_DATA_SPECIFY_WORKSPACE_LIMIT", CONV_NUMERIC_LITERAL, API_DNN}}, // 2 - {"cudnnIndicesType_t", {"hipdnnIndicesType_t", CONV_TYPE, API_DNN}}, {"CUDNN_32BIT_INDICES", {"HIPDNN_32BIT_INDICES", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_64BIT_INDICES", {"HIPDNN_64BIT_INDICES", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 {"CUDNN_16BIT_INDICES", {"HIPDNN_16BIT_INDICES", CONV_NUMERIC_LITERAL, API_DNN}}, // 2 {"CUDNN_8BIT_INDICES", {"HIPDNN_8BIT_INDICES", CONV_NUMERIC_LITERAL, API_DNN}}, // 3 - {"cudnnSoftmaxAlgorithm_t", {"hipdnnSoftmaxAlgorithm_t", CONV_TYPE, API_DNN}}, {"CUDNN_SOFTMAX_FAST", {"HIPDNN_SOFTMAX_FAST", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_SOFTMAX_ACCURATE", {"HIPDNN_SOFTMAX_ACCURATE", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 {"CUDNN_SOFTMAX_LOG", {"HIPDNN_SOFTMAX_LOG", CONV_NUMERIC_LITERAL, API_DNN}}, // 2 - {"cudnnSoftmaxMode_t", {"hipdnnSoftmaxMode_t", CONV_TYPE, API_DNN}}, {"CUDNN_SOFTMAX_MODE_INSTANCE", {"HIPDNN_SOFTMAX_MODE_INSTANCE", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_SOFTMAX_MODE_CHANNEL", {"HIPDNN_SOFTMAX_MODE_CHANNEL", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 - {"cudnnPoolingMode_t", {"hipdnnPoolingMode_t", CONV_TYPE, API_DNN}}, {"CUDNN_POOLING_MAX", {"HIPDNN_POOLING_MAX", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 {"CUDNN_POOLING_AVERAGE_COUNT_INCLUDE_PADDING", {"HIPDNN_POOLING_AVERAGE_COUNT_INCLUDE_PADDING", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 {"CUDNN_POOLING_AVERAGE_COUNT_EXCLUDE_PADDING", {"HIPDNN_POOLING_AVERAGE_COUNT_EXCLUDE_PADDING", CONV_NUMERIC_LITERAL, API_DNN}}, // 2 {"CUDNN_POOLING_MAX_DETERMINISTIC", {"HIPDNN_POOLING_MAX_DETERMINISTIC", CONV_NUMERIC_LITERAL, API_DNN}}, // 3 + {"cudnnActivationMode_t", {"hipdnnActivationMode_t", CONV_TYPE, API_DNN}}, + {"CUDNN_ACTIVATION_SIGMOID", {"HIPDNN_ACTIVATION_SIGMOID", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 + {"CUDNN_ACTIVATION_RELU", {"HIPDNN_ACTIVATION_RELU", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 + {"CUDNN_ACTIVATION_TANH", {"HIPDNN_ACTIVATION_TANH", CONV_NUMERIC_LITERAL, API_DNN}}, // 2 + {"CUDNN_ACTIVATION_CLIPPED_RELU", {"HIPDNN_ACTIVATION_CLIPPED_RELU", CONV_NUMERIC_LITERAL, API_DNN}}, // 3 + {"CUDNN_ACTIVATION_ELU", {"HIPDNN_ACTIVATION_ELU", CONV_NUMERIC_LITERAL, API_DNN}}, // 4 + {"CUDNN_ACTIVATION_IDENTITY", {"HIPDNN_ACTIVATION_PATHTRU", CONV_NUMERIC_LITERAL, API_DNN}}, // 5 + {"cudnnBatchNormMode_t", {"hipdnnBatchNormMode_t", CONV_TYPE, API_DNN}}, + {"CUDNN_BATCHNORM_PER_ACTIVATION", {"HIPDNN_BATCHNORM_PER_ACTIVATION", CONV_NUMERIC_LITERAL, API_DNN}}, // 0 + {"CUDNN_BATCHNORM_SPATIAL", {"HIPDNN_BATCHNORM_SPATIAL", CONV_NUMERIC_LITERAL, API_DNN}}, // 1 + {"CUDNN_BATCHNORM_SPATIAL_PERSISTENT", {"HIPDNN_BATCHNORM_SPATIAL_PERSISTENT", CONV_NUMERIC_LITERAL, API_DNN}}, // 2 + {"cudnnSamplerType_t", {"hipdnnSamplerType_t", CONV_TYPE, API_DNN, HIP_UNSUPPORTED}}, + {"CUDNN_SAMPLER_BILINEAR", {"HIPDNN_SAMPLER_BILINEAR", CONV_NUMERIC_LITERAL, API_DNN, HIP_UNSUPPORTED}}, // 0 {"cudnnGetVersion", {"hipdnnGetVersion", CONV_VERSION, API_DNN}}, {"cudnnGetCudartVersion", {"hipdnnGetCudartVersion", CONV_VERSION, API_DNN, HIP_UNSUPPORTED}}, @@ -3132,9 +3156,11 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cudnnCreateRNNDescriptor", {"hipdnnCreateRNNDescriptor", CONV_MATH_FUNC, API_DNN}}, {"cudnnSetStream", {"hipdnnSetStream", CONV_MATH_FUNC, API_DNN}}, {"cudnnGetStream", {"hipdnnGetStream", CONV_MATH_FUNC, API_DNN}}, + {"cudnnSetRNNDescriptor_v5", {"hipdnnSetRNNDescriptor_v5", CONV_MATH_FUNC, API_DNN}}, {"cudnnSetRNNDescriptor_v6", {"hipdnnSetRNNDescriptor_v6", CONV_MATH_FUNC, API_DNN}}, {"cudnnSetRNNDescriptor", {"hipdnnSetRNNDescriptor", CONV_MATH_FUNC, API_DNN}}, {"cudnnDropoutGetStatesSize", {"hipdnnDropoutGetStatesSize", CONV_MATH_FUNC, API_DNN}}, + {"cudnnDropoutGetReserveSpaceSize", {"hipdnnDropoutGetReserveSpaceSize", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, {"cudnnTransformTensor", {"hipdnnTransformTensor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, {"cudnnSetTensor4dDescriptor", {"hipdnnSetTensor4dDescriptor", CONV_MATH_FUNC, API_DNN}}, {"cudnnGetTensor4dDescriptor", {"hipdnnGetTensor4dDescriptor", CONV_MATH_FUNC, API_DNN}}, @@ -3158,16 +3184,25 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cudnnGetConvolutionBackwardDataAlgorithm", {"hipdnnGetConvolutionBackwardDataAlgorithm", CONV_MATH_FUNC, API_DNN}}, {"cudnnGetConvolutionBackwardDataAlgorithm_v7", {"hipdnnGetConvolutionBackwardDataAlgorithm_v7", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, {"cudnnGetConvolutionBackwardDataAlgorithmMaxCount", {"hipdnnGetConvolutionBackwardDataAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, - {"cudnnGetConvolutionForwardAlgorithmMaxCount", {"hipdnnGetConvolutionForwardAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetConvolutionForwardAlgorithmMaxCount", {"hipdnnGetConvolutionForwardAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, {"cudnnGetRNNLinLayerMatrixParams", {"hipdnnGetRNNLinLayerMatrixParams", CONV_MATH_FUNC, API_DNN}}, {"cudnnGetRNNLinLayerBiasParams", {"hipdnnGetRNNLinLayerBiasParams", CONV_MATH_FUNC, API_DNN}}, + {"cudnnSetRNNProjectionLayers", {"hipdnnSetRNNProjectionLayers", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetRNNProjectionLayers", {"hipdnnGetRNNProjectionLayers", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSetRNNAlgorithmDescriptor", {"hipdnnSetRNNAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetRNNDescriptor", {"hipdnnGetRNNDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, {"cudnnGetFilterNdDescriptor", {"hipdnnGetFilterNdDescriptor", CONV_MATH_FUNC, API_DNN}}, {"cudnnFindConvolutionBackwardDataAlgorithm", {"hipdnnFindConvolutionBackwardDataAlgorithm", CONV_MATH_FUNC, API_DNN}}, {"cudnnFindConvolutionBackwardDataAlgorithmEx", {"hipdnnFindConvolutionBackwardDataAlgorithmEx", CONV_MATH_FUNC, API_DNN}}, {"cudnnSetDropoutDescriptor", {"hipdnnSetDropoutDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnRestoreDropoutDescriptor", {"hipdnnRestoreDropoutDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetDropoutDescriptor", {"hipdnnGetDropoutDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetDropoutDescriptor", {"hipdnnGetDropoutDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, {"cudnnSetConvolution2dDescriptor", {"hipdnnSetConvolution2dDescriptor", CONV_MATH_FUNC, API_DNN}}, {"cudnnGetConvolution2dDescriptor", {"hipdnnGetConvolution2dDescriptor", CONV_MATH_FUNC, API_DNN}}, {"cudnnSetConvolutionMathType", {"hipdnnSetConvolutionMathType", CONV_MATH_FUNC, API_DNN}}, + {"cudnnDropoutForward", {"hipdnnDropoutForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnDropoutBackward", {"hipdnnDropoutBackward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, {"cudnnGetConvolutionMathType", {"hipdnnGetConvolutionMathType", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, {"cudnnSetConvolutionGroupCount", {"hipdnnSetConvolutionGroupCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, {"cudnnGetConvolutionGroupCount", {"hipdnnGetConvolutionGroupCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, @@ -3188,6 +3223,15 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cudnnRNNBackwardData", {"hipdnnRNNBackwardData", CONV_MATH_FUNC, API_DNN}}, {"cudnnSetFilter4dDescriptor", {"hipdnnSetFilter4dDescriptor", CONV_MATH_FUNC, API_DNN}}, {"cudnnSetRNNMatrixMathType", {"hipdnnSetRNNMatrixMathType", CONV_MATH_FUNC, API_DNN}}, + {"cudnnGetRNNMatrixMathType", {"hipdnnGetRNNMatrixMathType", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetRNNForwardInferenceAlgorithmMaxCount", {"hipdnnGetRNNForwardInferenceAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnFindRNNForwardInferenceAlgorithmEx", {"hipdnnFindRNNForwardInferenceAlgorithmEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetRNNForwardTrainingAlgorithmMaxCount", {"hipdnnGetRNNForwardTrainingAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnFindRNNForwardTrainingAlgorithmEx", {"hipdnnFindRNNForwardTrainingAlgorithmEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetRNNBackwardDataAlgorithmMaxCount", {"hipdnnGetRNNBackwardDataAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnFindRNNBackwardDataAlgorithmEx", {"hipdnnFindRNNBackwardDataAlgorithmEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetRNNBackwardWeightsAlgorithmMaxCount", {"hipdnnGetRNNBackwardWeightsAlgorithmMaxCount", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnFindRNNBackwardWeightsAlgorithmEx", {"hipdnnFindRNNBackwardWeightsAlgorithmEx", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, {"cudnnCreateConvolutionDescriptor", {"hipdnnCreateConvolutionDescriptor", CONV_MATH_FUNC, API_DNN}}, {"cudnnGetConvolutionForwardAlgorithm", {"hipdnnGetConvolutionForwardAlgorithm", CONV_MATH_FUNC, API_DNN}}, {"cudnnGetConvolutionForwardAlgorithm_v7", {"hipdnnGetConvolutionForwardAlgorithm_v7", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, @@ -3209,10 +3253,8 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cudnnDestroyReduceTensorDescriptor", {"hipdnnDestroyReduceTensorDescriptor", CONV_MATH_FUNC, API_DNN}}, {"cudnnDestroyPersistentRNNPlan", {"hipdnnDestroyPersistentRNNPlan", CONV_MATH_FUNC, API_DNN}}, {"cudnnDestroy", {"hipdnnDestroy", CONV_MATH_FUNC, API_DNN}}, - {"cudnnSoftmaxForward", {"hipdnnSoftmaxForward", CONV_MATH_FUNC, API_DNN}}, {"cudnnSoftmaxBackward", {"hipdnnSoftmaxBackward", CONV_MATH_FUNC, API_DNN}}, - {"cudnnCreatePoolingDescriptor", {"hipdnnCreatePoolingDescriptor", CONV_MATH_FUNC, API_DNN}}, {"cudnnSetPooling2dDescriptor", {"hipdnnSetPooling2dDescriptor", CONV_MATH_FUNC, API_DNN}}, {"cudnnGetPooling2dDescriptor", {"hipdnnGetPooling2dDescriptor", CONV_MATH_FUNC, API_DNN}}, @@ -3221,10 +3263,53 @@ const std::map CUDA_IDENTIFIER_MAP{ {"cudnnGetPoolingNdForwardOutputDim", {"hipdnnGetPoolingNdForwardOutputDim", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, {"cudnnGetPooling2dForwardOutputDim", {"hipdnnGetPooling2dForwardOutputDim", CONV_MATH_FUNC, API_DNN}}, {"cudnnDestroyPoolingDescriptor", {"hipdnnDestroyPoolingDescriptor", CONV_MATH_FUNC, API_DNN}}, - {"cudnnPoolingForward", {"hipdnnPoolingForward", CONV_MATH_FUNC, API_DNN}}, {"cudnnPoolingBackward", {"hipdnnPoolingBackward", CONV_MATH_FUNC, API_DNN}}, - + {"cudnnCreateActivationDescriptor", {"hipdnnCreateActivationDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnSetActivationDescriptor", {"hipdnnSetActivationDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnGetActivationDescriptor", {"hipdnnGetActivationDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnDestroyActivationDescriptor", {"hipdnnDestroyActivationDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnActivationForward", {"hipdnnActivationForward", CONV_MATH_FUNC, API_DNN}}, + {"cudnnActivationBackward", {"hipdnnActivationBackward", CONV_MATH_FUNC, API_DNN}}, + {"cudnnCreateLRNDescriptor", {"hipdnnCreateLRNDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnSetLRNDescriptor", {"hipdnnSetLRNDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnGetLRNDescriptor", {"hipdnnGetLRNDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnDestroyLRNDescriptor", {"hipdnnDestroyLRNDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnLRNCrossChannelForward", {"hipdnnLRNCrossChannelForward", CONV_MATH_FUNC, API_DNN}}, + {"cudnnLRNCrossChannelBackward", {"hipdnnLRNCrossChannelBackward", CONV_MATH_FUNC, API_DNN}}, + {"cudnnDivisiveNormalizationForward", {"hipdnnDivisiveNormalizationForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnDivisiveNormalizationBackward", {"hipdnnDivisiveNormalizationBackward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnDeriveBNTensorDescriptor", {"hipdnnDeriveBNTensorDescriptor", CONV_MATH_FUNC, API_DNN}}, + {"cudnnBatchNormalizationForwardTraining", {"hipdnnBatchNormalizationForwardTraining", CONV_MATH_FUNC, API_DNN}}, + {"cudnnBatchNormalizationForwardInference", {"hipdnnBatchNormalizationForwardInference", CONV_MATH_FUNC, API_DNN}}, + {"cudnnBatchNormalizationBackward", {"hipdnnBatchNormalizationBackward", CONV_MATH_FUNC, API_DNN}}, + {"cudnnCreateSpatialTransformerDescriptor", {"hipdnnCreateSpatialTransformerDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSetSpatialTransformerNdDescriptor", {"hipdnnSetSpatialTransformerNdDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnDestroySpatialTransformerDescriptor", {"hipdnnDestroySpatialTransformerDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSpatialTfGridGeneratorForward", {"hipdnnSpatialTfGridGeneratorForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSpatialTfGridGeneratorBackward", {"hipdnnSpatialTfGridGeneratorBackward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSpatialTfSamplerForward", {"hipdnnSpatialTfSamplerForward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSpatialTfSamplerBackward", {"hipdnnSpatialTfSamplerBackward", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnCreateCTCLossDescriptor", {"hipdnnCreateCTCLossDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSetCTCLossDescriptor", {"hipdnnSetCTCLossDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetCTCLossDescriptor", {"hipdnnGetCTCLossDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnDestroyCTCLossDescriptor", {"hipdnnDestroyCTCLossDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnCTCLoss", {"hipdnnCTCLoss", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetCTCLossWorkspaceSize", {"hipdnnGetCTCLossWorkspaceSize", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnCreateAlgorithmDescriptor", {"hipdnnCreateAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSetAlgorithmDescriptor", {"hipdnnSetAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetAlgorithmDescriptor", {"hipdnnGetAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnCopyAlgorithmDescriptor", {"hipdnnCopyAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnDestroyAlgorithmDescriptor", {"hipdnnDestroyAlgorithmDescriptor", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnCreateAlgorithmPerformance", {"hipdnnCreateAlgorithmPerformance", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSetAlgorithmPerformance", {"hipdnnSetAlgorithmPerformance", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetAlgorithmPerformance", {"hipdnnGetAlgorithmPerformance", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnDestroyAlgorithmPerformance", {"hipdnnDestroyAlgorithmPerformance", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetAlgorithmSpaceSize", {"hipdnnGetAlgorithmSpaceSize", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSaveAlgorithm", {"hipdnnSaveAlgorithm", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnRestoreAlgorithm", {"hipdnnRestoreAlgorithm", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnSetCallback", {"hipdnnSetCallback", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, + {"cudnnGetCallback", {"hipdnnGetCallback", CONV_MATH_FUNC, API_DNN, HIP_UNSUPPORTED}}, }; const std::map& CUDA_RENAMES_MAP() { From 40fb44dbe653e691e70d2962a6b684c685da8b97 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 21 May 2018 22:14:45 +0530 Subject: [PATCH 10/10] Fixed memcpy2D for pinned memory case using 2D kernel --- hipamd/src/hip_memory.cpp | 143 +++++++++++++++++++------------------- 1 file changed, 73 insertions(+), 70 deletions(-) diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp index e2202e7860..3abdaec25c 100644 --- a/hipamd/src/hip_memory.cpp +++ b/hipamd/src/hip_memory.cpp @@ -1208,7 +1208,6 @@ hipError_t hipMemcpyDtoD(hipDeviceptr_t dst, hipDeviceptr_t src, size_t sizeByte return ihipLogStatus(e); } - hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) { HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes); @@ -1227,7 +1226,6 @@ hipError_t hipMemcpyHtoH(void* dst, void* src, size_t sizeBytes) { return ihipLogStatus(e); } - hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind, hipStream_t stream) { HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, kind, stream); @@ -1235,7 +1233,6 @@ hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcp return ihipLogStatus(hip_internal::memcpyAsync(dst, src, sizeBytes, kind, stream)); } - hipError_t hipMemcpyHtoDAsync(hipDeviceptr_t dst, void* src, size_t sizeBytes, hipStream_t stream) { HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, stream); @@ -1258,51 +1255,6 @@ hipError_t hipMemcpyDtoHAsync(void* dst, hipDeviceptr_t src, size_t sizeBytes, h hip_internal::memcpyAsync(dst, src, sizeBytes, hipMemcpyDeviceToHost, stream)); } -// TODO - review and optimize -hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, - size_t height, hipMemcpyKind kind) { - if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return hipErrorInvalidValue; - - hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); - - hc::completion_future marker; - - hipError_t e = hipSuccess; - if((width == dpitch) && (width == spitch)) { - stream->locked_copySync((void*)dst, (void*)src, width*height, kind, false); - } else { - try { - for (int i = 0; i < height; ++i) { - stream->locked_copySync((unsigned char*)dst + i * dpitch, - (unsigned char*)src + i * spitch, width, kind); - } - } catch (ihipException& ex) { - e = ex._code; - } - } - - return e; -} - -hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, - size_t height, hipMemcpyKind kind) { - HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind); - hipError_t e = hipSuccess; - e = ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind); - return ihipLogStatus(e); -} - -hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { - HIP_INIT_SPECIAL_API((TRACE_MCMD), pCopy); - hipError_t e = hipSuccess; - if (pCopy == nullptr) { - e = hipErrorInvalidValue; - } - e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->widthInBytes, pCopy->srcHost, pCopy->srcPitch, - pCopy->widthInBytes, pCopy->height, hipMemcpyDefault); - return ihipLogStatus(e); -} - hipError_t hipMemcpy2DToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, wOffset, hOffset, src, spitch, width, height, kind); @@ -1534,16 +1486,15 @@ inline const T& clamp_integer(const T& x, const T& lower, const T& upper) { return std::min(upper, std::max(x, lower)); } -template -__global__ void hip_copy_n(T* dst, const T* src, size_t n) { - const uint32_t grid_dim = gridDim.x * blockDim.x; +template +__global__ void hip_copy2d_n(T* dst, const T* src, size_t width, size_t height, size_t destPitch, size_t srcPitch) { - size_t idx = blockIdx.x * block_dim + threadIdx.x; - while (idx < n) { - // __builtin_memcpy(reinterpret_cast(dst+idx), reinterpret_cast(src+idx), - // sizeof(T)); - dst[idx] = src[idx]; - idx += grid_dim; + size_t idx = blockIdx.x * blockDim.x + threadIdx.x; + size_t idy = blockIdx.y * blockDim.y + threadIdx.y; + if((idx < width) && (idy < height)){ + T *dstPtr = (T *)((uint8_t*) dst + idy * destPitch); + T *srcPtr = (T *)((uint8_t*) src + idy * srcPitch); + dstPtr[idx] = srcPtr[idx]; } } } // namespace @@ -1559,13 +1510,12 @@ void ihipMemsetKernel(hipStream_t stream, T* ptr, T val, size_t sizeBytes) { } template -void ihipMemcpyKernel(hipStream_t stream, T* dst, const T* src, size_t sizeBytes) { - static constexpr uint32_t block_dim_ = 256; - - const uint32_t grid_dim = clamp_integer(sizeBytes / block_dim_, 1, UINT32_MAX); - - hipLaunchKernelGGL(hip_copy_n, dim3(grid_dim), dim3{block_dim_}, 0u, stream, dst, src, - sizeBytes); +void ihipMemcpy2dKernel(hipStream_t stream, T* dst, const T* src, size_t width, size_t height, size_t destPitch, size_t srcPitch) { + size_t threadsPerBlock = 16; + uint32_t grid_dim_x = clamp_integer( (width+(threadsPerBlock-1)) / threadsPerBlock, 1, UINT32_MAX); + uint32_t grid_dim_y = clamp_integer( (height+(threadsPerBlock-1)) / threadsPerBlock, 1, UINT32_MAX); + hipLaunchKernelGGL(hip_copy2d_n, dim3(grid_dim_x,grid_dim_y), dim3(threadsPerBlock,threadsPerBlock), 0u, stream, dst, src, + width, height, destPitch, srcPitch); } typedef enum ihipMemsetDataType { @@ -1643,6 +1593,50 @@ int isLockedPointer(const void *ptr) return isLocked; } +// TODO - review and optimize +hipError_t ihipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, + size_t height, hipMemcpyKind kind) { + if (dst == nullptr || src == nullptr || width > dpitch || width > spitch) return hipErrorInvalidValue; + + hipStream_t stream = ihipSyncAndResolveStream(hipStreamNull); + int isLocked = 0; + if(kind == hipMemcpyHostToDevice) { + isLocked = isLockedPointer(src); + } else if(kind == hipMemcpyDeviceToHost) { + isLocked = isLockedPointer(dst); + } + + hc::completion_future marker; + + hipError_t e = hipSuccess; + if((width == dpitch) && (width == spitch)) { + stream->locked_copySync((void*)dst, (void*)src, width*height, kind, false); + } else { + try { + if(isLocked) { + for (int i = 0; i < height; ++i) + stream->locked_copySync((unsigned char*)dst + i * dpitch, + (unsigned char*)src + i * spitch, width, kind); + } else { + ihipMemcpy2dKernel (stream, static_cast (dst), static_cast (src), width, height, dpitch, spitch); + stream->locked_wait(); + } + } catch (ihipException& ex) { + e = ex._code; + } + } + + return e; +} + +hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, + size_t height, hipMemcpyKind kind) { + HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind); + hipError_t e = hipSuccess; + e = ihipMemcpy2D(dst, dpitch, src, spitch, width, height, kind); + return ihipLogStatus(e); +} + hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind, hipStream_t stream) { HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, dpitch, src, spitch, width, height, kind, stream); @@ -1658,14 +1652,12 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp hip_internal::memcpyAsync(dst, src, width*height, kind, stream); } else { try { - for (int i = 0; i < height; ++i) { - if(!isLocked) { + if(!isLocked){ + for (int i = 0; i < height; ++i) e = hip_internal::memcpyAsync((unsigned char*)dst + i * dpitch, (unsigned char*)src + i * spitch, width, kind, stream); - } else{ - size_t sizeBytes = width*height; - ihipMemcpyKernel (stream, static_cast (dst), static_cast (src), sizeBytes/sizeof(uint32_t)); - } + } else{ + ihipMemcpy2dKernel (stream, static_cast (dst), static_cast (src), width, height, dpitch, spitch); } } catch (ihipException& ex) { e = ex._code; @@ -1675,6 +1667,17 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp return ihipLogStatus(e); } +hipError_t hipMemcpyParam2D(const hip_Memcpy2D* pCopy) { + HIP_INIT_SPECIAL_API((TRACE_MCMD), pCopy); + hipError_t e = hipSuccess; + if (pCopy == nullptr) { + e = hipErrorInvalidValue; + } + e = ihipMemcpy2D(pCopy->dstArray->data, pCopy->widthInBytes, pCopy->srcHost, pCopy->srcPitch, + pCopy->widthInBytes, pCopy->height, hipMemcpyDefault); + return ihipLogStatus(e); +} + // TODO-sync: function is async unless target is pinned host memory - then these are fully sync. hipError_t hipMemsetAsync(void* dst, int value, size_t sizeBytes, hipStream_t stream) { HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, value, sizeBytes, stream);