Implement hipGetSymbolAddress and hipGetSymbolSize

[ROCm/hip commit: 73616582d6]

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`                                        |                                                |

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.

projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_functions.cpp

@@ -46,8 +46,8 @@ const std::map<llvm::StringRef, hipCounter> 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

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

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(

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<const char*>(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<const char*>(symbolName));
+}
+
+
 //---
 hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) {
     HIP_INIT_SPECIAL_API((TRACE_MCMD), dst, src, sizeBytes, kind);