diff --git a/src/hip_hcc.h b/src/hip_hcc.h
index 3074a4f121..c4b092aa5f 100644
--- a/src/hip_hcc.h
+++ b/src/hip_hcc.h
@@ -32,7 +32,7 @@ THE SOFTWARE.
 #error("This version of HIP requires a newer version of HCC.");
 #endif
 
-#define USE_DISPATCH_HSA_KERNEL 0
+#define USE_DISPATCH_HSA_KERNEL 1
 //
 
 
diff --git a/src/hip_module.cpp b/src/hip_module.cpp
index d58a476f77..badeac2dcd 100644
--- a/src/hip_module.cpp
+++ b/src/hip_module.cpp
@@ -163,7 +163,12 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){
     return ihipLogStatus(ret);
 }
 
-hipError_t hipModuleUnload(hipModule_t hmod){
+hipError_t hipModuleUnload(hipModule_t hmod)
+{
+    // TODO - improve this synchronization so it is thread-safe.
+    // Currently we want for all inflight activity to complete, but don't prevent another
+    // thread from launching new kernels before we finish this operation.
+    ihipSynchronize();
     hipError_t ret = hipSuccess;
     hsa_status_t status = hsa_executable_destroy(hmod->executable);
     if(status != HSA_STATUS_SUCCESS)