diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h
index ef591b8c51..a3a0963ba9 100644
--- a/projects/hip/include/hip/hcc_detail/hip_runtime.h
+++ b/projects/hip/include/hip/hcc_detail/hip_runtime.h
@@ -379,7 +379,7 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri
 #elif defined(__clang__) && defined(__HIP__)
 
 #define HIP_KERNEL_NAME(...) __VA_ARGS__
-#define HIP_SYMBOL(X) #X
+#define HIP_SYMBOL(X) X
 
 typedef int hipLaunchParm;
 
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 d35d426a5e..2b5649e721 100644
--- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h
+++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h
@@ -1488,18 +1488,18 @@ hipError_t hipMemcpyDtoDAsync(hipDeviceptr_t dst, hipDeviceptr_t src, size_t siz
 hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
     hipModule_t hmod, const char* name);
 
-hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName);
-hipError_t hipGetSymbolSize(size_t* size, const void* symbolName);
-hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src,
+hipError_t hipGetSymbolAddress(void** devPtr, const void* symbol);
+hipError_t hipGetSymbolSize(size_t* size, const void* symbol);
+hipError_t hipMemcpyToSymbol(const void* symbol, const void* src,
                              size_t sizeBytes, size_t offset __dparm(0),
                              hipMemcpyKind kind __dparm(hipMemcpyHostToDevice));
-hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src,
+hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src,
                                   size_t sizeBytes, size_t offset,
                                   hipMemcpyKind kind, hipStream_t stream __dparm(0));
-hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName,
+hipError_t hipMemcpyFromSymbol(void* dst, const void* symbol,
                                size_t sizeBytes, size_t offset __dparm(0),
                                hipMemcpyKind kind __dparm(hipMemcpyDeviceToHost));
-hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName,
+hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbol,
                                     size_t sizeBytes, size_t offset,
                                     hipMemcpyKind kind,
                                     hipStream_t stream __dparm(0));
@@ -3517,6 +3517,46 @@ static hipError_t __host__ inline hipOccupancyMaxActiveBlocksPerMultiprocessorWi
 }
 #endif  // defined(__cplusplus) && !defined(__HCC__) && defined(__clang__) && defined(__HIP__)
 
+#if defined(__cplusplus) && !defined(__HCC__)
+
+template <typename T>
+hipError_t hipGetSymbolAddress(void** devPtr, const T &symbol) {
+  return ::hipGetSymbolAddress(devPtr, (const void *)&symbol);
+}
+
+template <typename T>
+hipError_t hipGetSymbolSize(size_t* size, const T &symbol) {
+  return ::hipGetSymbolSize(size, (const void *)&symbol);
+}
+
+template <typename T>
+hipError_t hipMemcpyToSymbol(const T& symbol, const void* src, size_t sizeBytes,
+                             size_t offset __dparm(0),
+                             hipMemcpyKind kind __dparm(hipMemcpyHostToDevice)) {
+  return ::hipMemcpyToSymbol((const void*)&symbol, src, sizeBytes, offset, kind);
+}
+
+template <typename T>
+hipError_t hipMemcpyToSymbolAsync(const T& symbol, const void* src, size_t sizeBytes, size_t offset,
+                                  hipMemcpyKind kind, hipStream_t stream __dparm(0)) {
+  return ::hipMemcpyToSymbolAsync((const void*)&symbol, src, sizeBytes, offset, kind, stream);
+}
+
+template <typename T>
+hipError_t hipMemcpyFromSymbol(void* dst, const T &symbol,
+                               size_t sizeBytes, size_t offset __dparm(0),
+                               hipMemcpyKind kind __dparm(hipMemcpyDeviceToHost)) {
+  return ::hipMemcpyFromSymbol(dst, (const void*)&symbol, sizeBytes, offset, kind);
+}
+
+template <typename T>
+hipError_t hipMemcpyFromSymbolAsync(void* dst, const T& symbol, size_t sizeBytes, size_t offset,
+                                    hipMemcpyKind kind, hipStream_t stream __dparm(0)) {
+  return ::hipMemcpyFromSymbolAsync(dst, (const void*)&symbol, sizeBytes, offset, kind, stream);
+}
+
+#endif
+
 #if USE_PROF_API
 #include <hip/hcc_detail/hip_prof_str.h>
 #endif
diff --git a/projects/hip/vdi/hip_internal.hpp b/projects/hip/vdi/hip_internal.hpp
index d5d35fcf1d..92397d1255 100755
--- a/projects/hip/vdi/hip_internal.hpp
+++ b/projects/hip/vdi/hip_internal.hpp
@@ -237,6 +237,8 @@ private:
 
   std::unordered_map<const void*, DeviceFunction > functions_;
   std::unordered_multimap<std::string, DeviceVar > vars_;
+  // Map from the host shadow symbol to its device name.
+  std::unordered_map<const void*, std::string> symbols_;
 
   static PlatformState* platform_;
 
@@ -251,8 +253,10 @@ public:
   std::vector< std::pair<hipModule_t, bool> >* unregisterVar(hipModule_t hmod);
 
 
+  bool findSymbol(const void *hostVar, std::string &devName);
   PlatformState::DeviceVar* findVar(std::string hostVar, int deviceId, hipModule_t hmod);
-  void registerVar(const void* hostvar, const DeviceVar& var);
+  void registerVarSym(const void *hostVar, const char *symbolName);
+  void registerVar(const char* symbolName, const DeviceVar& var);
   void registerFunction(const void* hostFunction, const DeviceFunction& func);
 
   bool registerModFuncs(std::vector<std::string>& func_names, hipModule_t* module);
@@ -260,7 +264,7 @@ public:
   bool createFunc(hipFunction_t* hfunc, hipModule_t hmod, const char* name);
   hipFunction_t getFunc(const void* hostFunction, int deviceId);
   bool getFuncAttr(const void* hostFunction, hipFuncAttributes* func_attr);
-  bool getGlobalVar(const void* hostVar, int deviceId, hipModule_t hmod,
+  bool getGlobalVar(const char* hostVar, int deviceId, hipModule_t hmod,
                     hipDeviceptr_t* dev_ptr, size_t* size_ptr);
   bool getTexRef(const char* hostVar, hipModule_t hmod, textureReference** texRef);
 
diff --git a/projects/hip/vdi/hip_memory.cpp b/projects/hip/vdi/hip_memory.cpp
index 083b855fa9..78288a6a78 100644
--- a/projects/hip/vdi/hip_memory.cpp
+++ b/projects/hip/vdi/hip_memory.cpp
@@ -717,15 +717,19 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) {
 };
 
 
-hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t count,
+hipError_t hipMemcpyToSymbol(const void* symbol, const void* src, size_t count,
                              size_t offset, hipMemcpyKind kind) {
-  HIP_INIT_API(hipMemcpyToSymbol, symbolName, src, count, offset, kind);
+  HIP_INIT_API(hipMemcpyToSymbol, symbol, src, count, offset, kind);
 
   size_t sym_size = 0;
   hipDeviceptr_t device_ptr = nullptr;
 
+  std::string symbolName;
+  if (!PlatformState::instance().findSymbol(symbol, symbolName)) {
+    HIP_RETURN(hipErrorInvalidSymbol);
+  }
   /* Get address and size for the global symbol */
-  if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr,
+  if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), nullptr,
                                               &device_ptr, &sym_size)) {
     HIP_RETURN(hipErrorInvalidSymbol);
   }
@@ -741,15 +745,19 @@ hipError_t hipMemcpyToSymbol(const void* symbolName, const void* src, size_t cou
   HIP_RETURN(hipMemcpy(device_ptr, src, count, kind));
 }
 
-hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count,
+hipError_t hipMemcpyFromSymbol(void* dst, const void* symbol, size_t count,
                                size_t offset, hipMemcpyKind kind) {
-  HIP_INIT_API(hipMemcpyFromSymbol, symbolName, dst, count, offset, kind);
+  HIP_INIT_API(hipMemcpyFromSymbol, symbol, dst, count, offset, kind);
 
   size_t sym_size = 0;
   hipDeviceptr_t device_ptr = nullptr;
 
+  std::string symbolName;
+  if (!PlatformState::instance().findSymbol(symbol, symbolName)) {
+    HIP_RETURN(hipErrorInvalidSymbol);
+  }
   /* Get address and size for the global symbol */
-  if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr,
+  if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), nullptr,
                                               &device_ptr, &sym_size)) {
     HIP_RETURN(hipErrorInvalidSymbol);
   }
@@ -765,15 +773,19 @@ hipError_t hipMemcpyFromSymbol(void* dst, const void* symbolName, size_t count,
   HIP_RETURN(hipMemcpy(dst, device_ptr, count, kind));
 }
 
-hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, size_t count,
+hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t count,
                                   size_t offset, hipMemcpyKind kind, hipStream_t stream) {
-  HIP_INIT_API(hipMemcpyToSymbolAsync, symbolName, src, count, offset, kind, stream);
+  HIP_INIT_API(hipMemcpyToSymbolAsync, symbol, src, count, offset, kind, stream);
 
   size_t sym_size = 0;
   hipDeviceptr_t device_ptr = nullptr;
 
+  std::string symbolName;
+  if (!PlatformState::instance().findSymbol(symbol, symbolName)) {
+    HIP_RETURN(hipErrorInvalidSymbol);
+  }
   /* Get address and size for the global symbol */
-  if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr,
+  if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), nullptr,
                                               &device_ptr, &sym_size)) {
     HIP_RETURN(hipErrorInvalidSymbol);
   }
@@ -789,15 +801,19 @@ hipError_t hipMemcpyToSymbolAsync(const void* symbolName, const void* src, size_
   HIP_RETURN(hipMemcpyAsync(device_ptr, src, count, kind, stream));
 }
 
-hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbolName, size_t count,
+hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbol, size_t count,
                                     size_t offset, hipMemcpyKind kind, hipStream_t stream) {
-  HIP_INIT_API(hipMemcpyFromSymbolAsync, symbolName, dst, count, offset, kind, stream);
+  HIP_INIT_API(hipMemcpyFromSymbolAsync, symbol, dst, count, offset, kind, stream);
 
   size_t sym_size = 0;
   hipDeviceptr_t device_ptr = nullptr;
 
+  std::string symbolName;
+  if (!PlatformState::instance().findSymbol(symbol, symbolName)) {
+    HIP_RETURN(hipErrorInvalidSymbol);
+  }
   /* Get address and size for the global symbol */
-  if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr,
+  if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), nullptr,
                                               &device_ptr, &sym_size)) {
     HIP_RETURN(hipErrorInvalidSymbol);
   }
diff --git a/projects/hip/vdi/hip_platform.cpp b/projects/hip/vdi/hip_platform.cpp
index 0ddc2fe8af..ca79c5300d 100755
--- a/projects/hip/vdi/hip_platform.cpp
+++ b/projects/hip/vdi/hip_platform.cpp
@@ -266,7 +266,21 @@ PlatformState::DeviceVar* PlatformState::findVar(std::string hostVar, int device
   return dvar;
 }
 
-void PlatformState::registerVar(const void* hostvar,
+bool PlatformState::findSymbol(const void *hostVar, std::string &symbolName) {
+  auto it = symbols_.find(hostVar);
+  if (it != symbols_.end()) {
+    symbolName = it->second;
+    return true;
+  }
+  return false;
+}
+
+void PlatformState::registerVarSym(const void *hostVar, const char *symbolName) {
+  amd::ScopedLock lock(lock_);
+  symbols_.insert(std::make_pair(hostVar, std::string(symbolName)));
+}
+
+void PlatformState::registerVar(const char* hostvar,
                                 const DeviceVar& rvar) {
   amd::ScopedLock lock(lock_);
   vars_.insert(std::make_pair(std::string(reinterpret_cast<const char*>(hostvar)), rvar));
@@ -447,7 +461,7 @@ bool PlatformState::getTexRef(const char* hostVar, hipModule_t hmod, textureRefe
   return true;
 }
 
-bool PlatformState::getGlobalVar(const void* hostVar, int deviceId, hipModule_t hmod,
+bool PlatformState::getGlobalVar(const char* hostVar, int deviceId, hipModule_t hmod,
                                  hipDeviceptr_t* dev_ptr, size_t* size_ptr) {
   amd::ScopedLock lock(lock_);
   DeviceVar* dvar = findVar(std::string(reinterpret_cast<const char*>(hostVar)), deviceId, hmod);
@@ -542,6 +556,7 @@ extern "C" void __hipRegisterVar(
     std::vector<PlatformState::RegisteredVar>{g_devices.size()}, false };
 
   PlatformState::instance().registerVar(hostVar, dvar);
+  PlatformState::instance().registerVarSym(var, deviceVar);
 }
 
 extern "C" void __hipUnregisterFatBinary(std::vector< std::pair<hipModule_t, bool> >* modules)
@@ -643,18 +658,30 @@ extern "C" hipError_t hipLaunchByPtr(const void *hostFunction)
     exec.sharedMem_, exec.hStream_, nullptr, extra));
 }
 
-hipError_t hipGetSymbolAddress(void** devPtr, const void* symbolName) {
+hipError_t hipGetSymbolAddress(void** devPtr, const void* symbol) {
+  HIP_INIT_API(hipGetSymbolAddress, devPtr, symbol);
+
+  std::string symbolName;
+  if (!PlatformState::instance().findSymbol(symbol, symbolName)) {
+    HIP_RETURN(hipErrorInvalidSymbol);
+  }
   size_t size = 0;
-  if(!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr,
+  if(!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), nullptr,
                                              devPtr, &size)) {
     HIP_RETURN(hipErrorInvalidSymbol);
   }
   HIP_RETURN(hipSuccess);
 }
 
-hipError_t hipGetSymbolSize(size_t* sizePtr, const void* symbolName) {
+hipError_t hipGetSymbolSize(size_t* sizePtr, const void* symbol) {
+  HIP_INIT_API(hipGetSymbolSize, sizePtr, symbol);
+
+  std::string symbolName;
+  if (!PlatformState::instance().findSymbol(symbol, symbolName)) {
+    HIP_RETURN(hipErrorInvalidSymbol);
+  }
   hipDeviceptr_t devPtr = nullptr;
-  if (!PlatformState::instance().getGlobalVar(symbolName, ihipGetDevice(), nullptr,
+  if (!PlatformState::instance().getGlobalVar(symbolName.c_str(), ihipGetDevice(), nullptr,
                                               &devPtr, sizePtr)) {
     HIP_RETURN(hipErrorInvalidSymbol);
   }