diff --git a/samples/common/hsa_base_util.cpp b/samples/common/hsa_base_util.cpp
index bd5a5cf9e6..3d35093764 100644
--- a/samples/common/hsa_base_util.cpp
+++ b/samples/common/hsa_base_util.cpp
@@ -1,4 +1,5 @@
 #include "hsa_base_util.h"
+#include "HSAILAmdExt.h"
 
 
 void HSA_UTIL::GetHsailNameAndKernelName(char * file_name, char *kernel_name)
@@ -55,6 +56,7 @@ bool HSA_UTIL::HsaInit()
 	err = hsa_queue_create(device, queue_size, HSA_QUEUE_TYPE_MULTI, NULL, NULL, 0, 0, &command_queue);
 	check(Creating the queue, err);
 
+        amd::hsail::registerExtensions();
         if (!tool.assembleFromFile(hail_file_name)) {
           std::cout << tool.output();
           return false;
diff --git a/samples/common/hsa_rsrc_factory.cpp b/samples/common/hsa_rsrc_factory.cpp
index d1bac57c77..81eeaa5a7b 100755
--- a/samples/common/hsa_rsrc_factory.cpp
+++ b/samples/common/hsa_rsrc_factory.cpp
@@ -11,6 +11,7 @@
 #include "hsa.h"
 #include "hsa_rsrc_factory.hpp"
 #include "hsa_ext_finalize.h"
+#include "HSAILAmdExt.h"
 
 #include "common.hpp"
 
@@ -317,6 +318,8 @@ bool HsaRsrcFactory::LoadAndFinalize(AgentInfo *agent_info,
   check("Error in Finding the Symbol Offset for the Kernel", status);
   */
 
+  amd::hsail::registerExtensions();
+
   // Copy handle of Brig object
   hsa_ext_module_t brig_module_v3;
   if (!tool.assembleFromFile(brig_path)) {