From bc455ccf502eb68fe587ea49156b393072c34f2b Mon Sep 17 00:00:00 2001 From: Michael Kuron Date: Sun, 4 Nov 2018 10:39:34 +0100 Subject: [PATCH] Implement hipGetSymbolAddress and hipGetSymbolSize [ROCm/hip commit: 73616582d6f8858d141933858f4b6db5d7cd51c5] --- ..._Runtime_API_functions_supported_by_HIP.md | 8 +++--- .../hip/docs/markdown/hip_kernel_language.md | 2 +- .../src/CUDA2HIP_Runtime_API_functions.cpp | 4 +-- .../include/hip/hcc_detail/hip_runtime_api.h | 26 +++++++++++++++++++ .../include/hip/nvcc_detail/hip_runtime_api.h | 8 ++++++ projects/hip/src/hip_memory.cpp | 17 ++++++++++++ 6 files changed, 58 insertions(+), 7 deletions(-) diff --git a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index 1a55667f82..7e5cd6fc3d 100644 --- a/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/projects/hip/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md @@ -122,8 +122,8 @@ | `cudaFreeHost` | `hipHostFree` | | `cudaFreeMipmappedArray` | | | `cudaGetMipmappedArrayLevel` | | -| `cudaGetSymbolAddress` | | -| `cudaGetSymbolSize` | | +| `cudaGetSymbolAddress` | `hipGetSymbolAddress` | +| `cudaGetSymbolSize` | `hipGetSymbolSize` | | `cudaHostAlloc` | `hipHostMalloc` | | `cudaHostGetDevicePointer` | `hipHostGetDevicePointer` | | `cudaHostGetFlags` | `hipHostGetFlags` | @@ -373,8 +373,8 @@ | `cudaCreateChannelDesc` | `hipCreateChannelDesc` | | `cudaFuncGetAttributes` | | | `cudaFuncSetCacheConfig` | | -| `cudaGetSymbolAddress` | | -| `cudaGetSymbolSize` | | +| `cudaGetSymbolAddress` | `hipGetSymbolAddress` | +| `cudaGetSymbolSize` | `hipGetSymbolSize` | | `cudaGetTextureAlignmentOffset` | | | `cudaLaunch` | | | `cudaLaunchKernel` | | diff --git a/projects/hip/docs/markdown/hip_kernel_language.md b/projects/hip/docs/markdown/hip_kernel_language.md index 7c209acadf..d69f5a04a8 100644 --- a/projects/hip/docs/markdown/hip_kernel_language.md +++ b/projects/hip/docs/markdown/hip_kernel_language.md @@ -159,7 +159,7 @@ void callMyKernel() ## Variable-Type Qualifiers ### `__constant__` -The `__constant__` keyword is supported. The host writes constant memory before launching the kernel; from the GPU, this memory is read-only during kernel execution. The functions for accessing constant memory (hipGetSymbolAddress(), hipGetSymbolSize(), hipMemcpyToSymbol(), hipMemcpyToSymbolAsync, hipMemcpyFromSymbol, hipMemcpyFromSymbolAsync) are under development. +The `__constant__` keyword is supported. The host writes constant memory before launching the kernel; from the GPU, this memory is read-only during kernel execution. The functions for accessing constant memory (hipGetSymbolAddress(), hipGetSymbolSize(), hipMemcpyToSymbol(), hipMemcpyToSymbolAsync(), hipMemcpyFromSymbol(), hipMemcpyFromSymbolAsync()) are available. ### `__shared__` The `__shared__` keyword is supported. diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp index 6c81de2817..12427b2ff4 100644 --- a/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp +++ b/projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp @@ -46,8 +46,8 @@ const std::map CUDA_RUNTIME_FUNCTION_MAP{ {"cudaArrayGetInfo", {"hipArrayGetInfo", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}}, {"cudaFreeMipmappedArray", {"hipFreeMipmappedArray", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}}, {"cudaGetMipmappedArrayLevel", {"hipGetMipmappedArrayLevel", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}}, - {"cudaGetSymbolAddress", {"hipGetSymbolAddress", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}}, - {"cudaGetSymbolSize", {"hipGetSymbolSize", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}}, + {"cudaGetSymbolAddress", {"hipGetSymbolAddress", CONV_MEMORY, API_RUNTIME}}, + {"cudaGetSymbolSize", {"hipGetSymbolSize", CONV_MEMORY, API_RUNTIME}}, {"cudaMemPrefetchAsync", {"hipMemPrefetchAsync", CONV_MEMORY, API_RUNTIME, HIP_UNSUPPORTED}}, // // API_Driver ANALOGUE (cuMemPrefetchAsync) // malloc diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index c9ff32d197..8b8bbe9e4f 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -1385,6 +1385,32 @@ hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t siz size_t offset __dparm(0), hipMemcpyKind kind __dparm(hipMemcpyHostToDevice)); +/** + * @brief Copies the memory address of symbol @p symbolName to @p devPtr + * + * @param[in] symbolName - Symbol on device + * @param[out] devPtr - Pointer to a pointer to the memory referred to by the symbol + * @return #... + * + * @see hipGetSymbolSize, hipMemcpyToSymbol, hipMemcpyFromSymbol, hipMemcpyToSymbolAsync, + * hipMemcpyFromSymbolAsync + */ +hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName); + + +/** + * @brief Copies the size of symbol @p symbolName to @p size + * + * @param[in] symbolName - Symbol on device + * @param[out] size - Pointer to the size of the symbol + * @return #... + * + * @see hipGetSymbolSize, hipMemcpyToSymbol, hipMemcpyFromSymbol, hipMemcpyToSymbolAsync, + * hipMemcpyFromSymbolAsync + */ +hipError_t hipGetSymbolSize(size_t* size, const void* symbolName); + + /** * @brief Copies @p sizeBytes bytes from the memory area pointed to by @p src to the memory area * pointed to by @p offset bytes from the start of symbol @p symbol diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index e68e8ac328..88867567c9 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -551,6 +551,14 @@ inline static hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolN dst, symbolName, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(kind), stream)); } +inline static hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) { + return hipCUDAErrorTohipError(cudaGetSymbolAddress(devPtr, symbolName); +} + +inline static hipError_t hipGetSymbolSize(size_t* size, const void* symbolName) { + return hipCUDAErrorTohipError(cudaGetSymbolSize(size, symbolName); +} + inline static hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, size_t height, hipMemcpyKind kind) { return hipCUDAErrorTohipError( diff --git a/projects/hip/src/hip_memory.cpp b/projects/hip/src/hip_memory.cpp index 4ea5b24f43..87e5f97e73 100644 --- a/projects/hip/src/hip_memory.cpp +++ b/projects/hip/src/hip_memory.cpp @@ -1099,6 +1099,23 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co return ihipLogStatus(e); } + +hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) { + HIP_INIT_SPECIAL_API((TRACE_MCMD), devPtr, symbolName); + + size_t size = 0; + return hipModuleGetGlobal(devPtr, &size, 0, static_cast(symbolName)); +} + + +hipError_t hipGetSymbolSize(size_t* size, const void* symbolName) { + HIP_INIT_SPECIAL_API((TRACE_MCMD), size, symbolName); + + void* devPtr = nullptr; + return hipModuleGetGlobal(&devPtr, size, 0, static_cast(symbolName)); +} + + //--- hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) { HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, kind);