Enabled USE_DISPATCH_HSA_KERNEL, with serialization in hipModuleUnload.

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

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)