diff --git a/api/hip/hip_internal.hpp b/api/hip/hip_internal.hpp
index 1f7e31c9aa..0df562065c 100644
--- a/api/hip/hip_internal.hpp
+++ b/api/hip/hip_internal.hpp
@@ -175,6 +175,8 @@ extern hipError_t ihipDeviceGetCount(int* count);
 extern int ihipGetDevice();
 extern hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags);
 extern amd::Memory* getMemoryObject(const void* ptr, size_t& offset);
+extern bool CL_CALLBACK getSvarInfo(cl_program program, std::string var_name, void** var_addr,
+                                    size_t* var_size);
 
 #define HIP_RETURN(ret)          \
         hip::g_lastError = ret;  \
diff --git a/api/hip/hip_module.cpp b/api/hip/hip_module.cpp
index 77a8a34659..a66c2da9e9 100644
--- a/api/hip/hip_module.cpp
+++ b/api/hip/hip_module.cpp
@@ -143,6 +143,8 @@ hipError_t ihipModuleLoadData(hipModule_t *module, const void *image)
     return hipErrorOutOfMemory;
   }
 
+  program->setVarInfoCallBack(&getSvarInfo);
+
   if (CL_SUCCESS != program->addDeviceProgram(*hip::getCurrentContext()->devices()[0], image, ElfSize(image)) ||
     CL_SUCCESS != program->build(hip::getCurrentContext()->devices(), nullptr, nullptr, nullptr)) {
     return hipErrorUnknown;