From a38366c81e391361f4559f06c0cf091905e70b33 Mon Sep 17 00:00:00 2001 From: Michael Kuron Date: Sun, 4 Nov 2018 10:39:34 +0100 Subject: [PATCH 1/7] Implement hipGetSymbolAddress and hipGetSymbolSize [ROCm/clr commit: 0b6f5791f82a8a0ef714cbd964a74ede9bbe1c4d] --- ..._Runtime_API_functions_supported_by_HIP.md | 8 +++--- .../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/clr/hipamd/src/hip_memory.cpp | 17 ++++++++++++ 6 files changed, 58 insertions(+), 7 deletions(-) diff --git a/projects/clr/hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md b/projects/clr/hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md index 1a55667f82..7e5cd6fc3d 100644 --- a/projects/clr/hipamd/docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md +++ b/projects/clr/hipamd/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/clr/hipamd/docs/markdown/hip_kernel_language.md b/projects/clr/hipamd/docs/markdown/hip_kernel_language.md index 7c209acadf..d69f5a04a8 100644 --- a/projects/clr/hipamd/docs/markdown/hip_kernel_language.md +++ b/projects/clr/hipamd/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/clr/hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp index 6c81de2817..12427b2ff4 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp +++ b/projects/clr/hipamd/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/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index c9ff32d197..8b8bbe9e4f 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/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/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index e68e8ac328..88867567c9 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/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/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 4ea5b24f43..87e5f97e73 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/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); From 15353a7b266b80ebf0a9759ed02046306f191fee Mon Sep 17 00:00:00 2001 From: Michael Kuron Date: Sun, 4 Nov 2018 11:47:17 +0100 Subject: [PATCH 2/7] Document return values of hipMemcpyToSymbol, hipGetSymbolAddress [ROCm/clr commit: 8049cdafab3836eba831e30d39ce23376295177d] --- projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 8b8bbe9e4f..cfecb144c3 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -1390,7 +1390,7 @@ hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t siz * * @param[in] symbolName - Symbol on device * @param[out] devPtr - Pointer to a pointer to the memory referred to by the symbol - * @return #... + * @return #hipSuccess, #hipErrorNotInitialized, #hipErrorNotFound * * @see hipGetSymbolSize, hipMemcpyToSymbol, hipMemcpyFromSymbol, hipMemcpyToSymbolAsync, * hipMemcpyFromSymbolAsync @@ -1403,7 +1403,7 @@ hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName); * * @param[in] symbolName - Symbol on device * @param[out] size - Pointer to the size of the symbol - * @return #... + * @return #hipSuccess, #hipErrorNotInitialized, #hipErrorNotFound * * @see hipGetSymbolSize, hipMemcpyToSymbol, hipMemcpyFromSymbol, hipMemcpyToSymbolAsync, * hipMemcpyFromSymbolAsync From a857354072073d92b3aa3a8b98e8f73d0d1a33e8 Mon Sep 17 00:00:00 2001 From: Michael Kuron Date: Tue, 6 Nov 2018 09:54:34 +0100 Subject: [PATCH 3/7] Introduce ihipModuleGetGlobal [ROCm/clr commit: 4da2d92281bfd255c106022217c478ca8d0867f6] --- .../hipamd/include/hip/hcc_detail/hip_runtime_api.h | 3 +++ projects/clr/hipamd/src/hip_memory.cpp | 6 +++--- projects/clr/hipamd/src/hip_module.cpp | 11 ++++++++--- 3 files changed, 14 insertions(+), 6 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index cfecb144c3..3567a67854 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -2360,6 +2360,9 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func); hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod, const char* name); +hipError_t ihipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod, + const char* name); + hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name); /** * @brief builds module from code object which resides in host memory. Image is pointer to that diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 87e5f97e73..4322ada199 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -953,7 +953,7 @@ inline hipDeviceptr_t agent_address_for_symbol(const char* symbolName) { #if __hcc_workweek__ >= 17481 size_t byte_cnt = 0u; - hipModuleGetGlobal(&r, &byte_cnt, 0, symbolName); + ihipModuleGetGlobal(&r, &byte_cnt, 0, symbolName); #else auto ctx = ihipGetTlsDefaultCtx(); auto acc = ctx->getDevice()->_acc; @@ -1104,7 +1104,7 @@ 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)); + return ihipModuleGetGlobal(devPtr, &size, 0, static_cast(symbolName)); } @@ -1112,7 +1112,7 @@ 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)); + return ihipModuleGetGlobal(&devPtr, size, 0, static_cast(symbolName)); } diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index a6d486b6de..efb091f68a 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -463,14 +463,19 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t h const char* name) { HIP_INIT_API(dptr, bytes, hmod, name); - if (!dptr || !bytes) return ihipLogStatus(hipErrorInvalidValue); + return ihipLogStatus(ihipModuleGetGlobal(dptr, bytes, hmod, name)); +} - if (!name) return ihipLogStatus(hipErrorNotInitialized); +hipError_t ihipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t hmod, + const char* name) { + if (!dptr || !bytes) return hipErrorInvalidValue; + + if (!name) return hipErrorNotInitialized; const auto r = hmod ? read_agent_global_from_module(dptr, bytes, hmod, name) : read_agent_global_from_process(dptr, bytes, name); - return ihipLogStatus(r); + return r; } namespace From 35227acd6d8cb158a0b3437527a0735737074a68 Mon Sep 17 00:00:00 2001 From: Michael Kuron Date: Tue, 6 Nov 2018 11:39:34 +0100 Subject: [PATCH 4/7] Test for hipGetSymbolSize and hipGetSymbolAddress [ROCm/clr commit: fe2281f101c39d096e658a8598274491f5f41746] --- .../src/deviceLib/hipTestDeviceSymbol.cpp | 20 +++++++++++++++++++ 1 file changed, 20 insertions(+) diff --git a/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp b/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp index 4bac9a902b..a17dd75a31 100644 --- a/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp +++ b/projects/clr/hipamd/tests/src/deviceLib/hipTestDeviceSymbol.cpp @@ -40,6 +40,12 @@ __global__ void Assign(int* Out) { globalOut[tid] = globalIn[tid]; } +__device__ __constant__ int globalConst[NUM]; + +__global__ void checkAddress(int* addr, bool* out) { + *out = (globalConst == addr); +} + int main() { int *A, *Am, *B, *Ad, *C, *Cm; A = new int[NUM]; @@ -101,6 +107,20 @@ int main() { assert(A[i] == B[i]); assert(A[i] == C[i]); } + + bool *checkOkD; + bool checkOk = false; + size_t symbolSize = 0; + int *symbolAddress; + hipGetSymbolSize(&symbolSize, HIP_SYMBOL(globalConst)); + hipGetSymbolAddress((void**) &symbolAddress, HIP_SYMBOL(globalConst)); + hipMalloc((void**)&checkOkD, sizeof(bool)); + hipLaunchKernelGGL(checkAddress, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, symbolAddress, checkOkD); + hipMemcpy(&checkOk, checkOkD, sizeof(bool), hipMemcpyDeviceToHost); + hipFree(checkOkD); + assert(checkOk); + assert(symbolSize == SIZE); + hipHostFree(Am); hipHostFree(Cm); hipFree(Ad); From f8b1d20fd9cf5eef64312159433e96ac4e1e594e Mon Sep 17 00:00:00 2001 From: Michael Kuron Date: Tue, 6 Nov 2018 12:02:21 +0100 Subject: [PATCH 5/7] hipify-perl: add hipGetSymbolAddress and hipGetSymbolSize [ROCm/clr commit: 538a8939a6b184d42ff65aaf8cc8f0c6265d7272] --- projects/clr/hipamd/bin/hipify-perl | 3 +++ 1 file changed, 3 insertions(+) diff --git a/projects/clr/hipamd/bin/hipify-perl b/projects/clr/hipamd/bin/hipify-perl index 3ea201d14f..792bf86d94 100755 --- a/projects/clr/hipamd/bin/hipify-perl +++ b/projects/clr/hipamd/bin/hipify-perl @@ -287,6 +287,9 @@ while (@ARGV) { $ft{'mem'} += s/\bcudaMemcpy2DToArray\b/hipMemcpy2DToArray/g; $ft{'mem'} += s/\bcudaMemcpyToArray\b/hipMemcpyToArray/g; + $ft{'mem'} += s/\bcudaGetSymbolAddress\b/hipGetSymbolAddress/g; + $ft{'mem'} += s/\bcudaGetSymbolSize\b/hipGetSymbolSize/g; + #-------- # Memory management: $ft{'mem'} += s/\bcudaMalloc\b/hipMalloc/g; From a0c35ab823e1d4a625206557ba7672a9af65bc1f Mon Sep 17 00:00:00 2001 From: Michael Kuron Date: Tue, 6 Nov 2018 20:46:30 +0100 Subject: [PATCH 6/7] Use correct trace macro in hipGetSymbolAddress/hipGetSymbolSize [ROCm/clr commit: 357dc8be117f8714bfd538c1a8fac9506af35e00] --- projects/clr/hipamd/src/hip_memory.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 4322ada199..d02bb5acb8 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -1101,7 +1101,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t co hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) { - HIP_INIT_SPECIAL_API((TRACE_MCMD), devPtr, symbolName); + HIP_INIT_API(devPtr, symbolName); size_t size = 0; return ihipModuleGetGlobal(devPtr, &size, 0, static_cast(symbolName)); @@ -1109,7 +1109,7 @@ hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) { hipError_t hipGetSymbolSize(size_t* size, const void* symbolName) { - HIP_INIT_SPECIAL_API((TRACE_MCMD), size, symbolName); + HIP_INIT_API(size, symbolName); void* devPtr = nullptr; return ihipModuleGetGlobal(&devPtr, size, 0, static_cast(symbolName)); From 4a646eed6c6f6485ce8c317ebfdd13cd8d830d48 Mon Sep 17 00:00:00 2001 From: Michael Kuron Date: Thu, 15 Nov 2018 09:48:00 +0100 Subject: [PATCH 7/7] Fix hipGetSymbolAddress/hipGetSymbolSize on nvcc [ROCm/clr commit: f71b8cbc7dac8a4c332df211a3b9bff1d4b05683] --- projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index 88867567c9..02c4b7ee61 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -552,11 +552,11 @@ inline static hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolN } inline static hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) { - return hipCUDAErrorTohipError(cudaGetSymbolAddress(devPtr, symbolName); + return hipCUDAErrorTohipError(cudaGetSymbolAddress(devPtr, symbolName)); } inline static hipError_t hipGetSymbolSize(size_t* size, const void* symbolName) { - return hipCUDAErrorTohipError(cudaGetSymbolSize(size, symbolName); + return hipCUDAErrorTohipError(cudaGetSymbolSize(size, symbolName)); } inline static hipError_t hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch,