Add additional controls for forcing serialization and blocking.

Move HIP_COHERENT_HOST_ALLOC so it is read once at init time.
Add HIP_LAUNCH_BLOCKING_KERNELS, HIP_API_BLOCKING.
Update docs on debug and chicken bits.

Conflicts:
	src/hip_hcc.cpp

src/hip_hcc.cpp

@@ -61,6 +61,9 @@ const char *API_COLOR = KGRN;
 const char *API_COLOR_END = KNRM;
 
 int HIP_LAUNCH_BLOCKING = 0;
+std::string HIP_LAUNCH_BLOCKING_KERNELS;
+std::vector<std::string> g_hipLaunchBlockingKernels;
+int HIP_API_BLOCKING = 0;
 
 int HIP_PRINT_ENV = 0;
 int HIP_TRACE_API= 0;
@@ -81,6 +84,8 @@ int HIP_DENY_PEER_ACCESS = 0;
 // Force async copies to actually use the synchronous copy interface.
 int HIP_FORCE_SYNC_COPY = 0;
 
+int HIP_COHERENT_HOST_ALLOC = 0;
+
 
 
 
@@ -358,13 +363,24 @@ LockedAccessor_StreamCrit_t ihipStream_t::lockopen_preKernelCommand()
 
 //---
 // Must be called after kernel finishes, this releases the lock on the stream so other commands can submit.
-void ihipStream_t::lockclose_postKernelCommand(hc::accelerator_view *av)
+void ihipStream_t::lockclose_postKernelCommand(const char * kernelName, hc::accelerator_view *av)
 {
+    bool blockThisKernel = false;
 
-    if (HIP_LAUNCH_BLOCKING) {
+    if (!g_hipLaunchBlockingKernels.empty()) {
+        std::string kernelNameString(kernelName);
+        for (auto o=g_hipLaunchBlockingKernels.begin(); o!=g_hipLaunchBlockingKernels.end(); o++) {
+            if ((*o == kernelNameString)) {
+                //printf ("force blocking for kernel %s\n", o->c_str());
+                blockThisKernel = true;
+            }
+        }
+    }
+
+    if (HIP_LAUNCH_BLOCKING || blockThisKernel) {
         // TODO - fix this so it goes through proper stream::wait() call.// direct wait OK since we know the stream is locked.
         av->wait(hc::hcWaitModeActive);
-        tprintf(DB_SYNC, " %s LAUNCH_BLOCKING for kernel completion\n", ToString(this).c_str());
+        tprintf(DB_SYNC, "%s LAUNCH_BLOCKING for kernel '%s' completion\n", ToString(this).c_str(), kernelName);
     }
 
     _criticalData.unlock(); // paired with lock from lockopen_preKernelCommand.
@@ -1243,7 +1259,15 @@ void ihipInit()
     //-- READ HIP_PRINT_ENV env first, since it has impact on later env var reading
 
     // TODO: In HIP/hcc, this variable blocks after both kernel commmands and data transfer.  Maybe should be bit-mask for each command type?
-    READ_ENV_I(release, HIP_LAUNCH_BLOCKING, CUDA_LAUNCH_BLOCKING, "Make HIP APIs 'host-synchronous', so they block until any kernel launches or data copy commands complete. Alias: CUDA_LAUNCH_BLOCKING." );
+    READ_ENV_I(release, HIP_LAUNCH_BLOCKING, CUDA_LAUNCH_BLOCKING, "Make HIP kernel launches 'host-synchronous', so they block until any kernel launches. Alias: CUDA_LAUNCH_BLOCKING." );
+    READ_ENV_S(release, HIP_LAUNCH_BLOCKING_KERNELS, 0, "Comma-separated list of kernel names to make host-synchronous, so they block until completed.");
+    if (!HIP_LAUNCH_BLOCKING_KERNELS.empty()) {
+        tokenize(HIP_LAUNCH_BLOCKING_KERNELS, ',', &g_hipLaunchBlockingKernels);
+    }
+    READ_ENV_I(release, HIP_API_BLOCKING, 0, "Make HIP APIs 'host-synchronous', so they block until completed.  Impacts hipMemcpyAsync, hipMemsetAsync." );
+    
+
+
     READ_ENV_C(release, HIP_DB, 0,  "Print debug info.  Bitmask (HIP_DB=0xff) or flags separated by '+' (HIP_DB=api+sync+mem+copy)", HIP_DB_callback);
     if ((HIP_DB & (1<<DB_API))  && (HIP_TRACE_API == 0)) {
         // Set HIP_TRACE_API default before we read it, so it is printed correctly.
@@ -1252,6 +1276,8 @@ void ihipInit()
 
 
 
+
+
     READ_ENV_I(release, HIP_TRACE_API, 0,  "Trace each HIP API call.  Print function name and return code to stderr as program executes.");
     READ_ENV_S(release, HIP_TRACE_API_COLOR, 0,  "Color to use for HIP_API.  None/Red/Green/Yellow/Blue/Magenta/Cyan/White");
     READ_ENV_I(release, HIP_PROFILE_API, 0,  "Add HIP API markers to ATP file generated with CodeXL. 0x1=short API name, 0x2=full API name including args.");
@@ -1262,10 +1288,14 @@ void ihipInit()
 
 
     READ_ENV_I(release, HIP_WAIT_MODE, 0, "Force synchronization mode. 1= force yield, 2=force spin, 0=defaults specified in application");
-    READ_ENV_I(release, HIP_FORCE_P2P_HOST, 0, "Force use of host/staging copy for peer-to-peer copies.1=always use copies, 2=always return false for hipDeviceCanAccessPeer");
-    READ_ENV_I(release, HIP_FORCE_SYNC_COPY, 0, "Force all copies (even hipMemcpyAsync) to use sync copies");
+    READ_ENV_I(release, HIP_FORCE_P2P_HOST, 0, "Force use of host/staging copy for peer-to-peer copies.1=always use copies, 2=always return false for hipDeviceCanAccessPeer"); 
+    READ_ENV_I(release, HIP_FORCE_SYNC_COPY, 0, "Force all copies (even hipMemcpyAsync) to use sync copies"); 
+
+    // TODO - review, can we remove this?
     READ_ENV_I(release, HIP_NUM_KERNELS_INFLIGHT, 128, "Max number of inflight kernels per stream before active synchronization is forced.");
 
+    READ_ENV_I(release, HIP_COHERENT_HOST_ALLOC, 0, "If set, all host memory will be allocated as fine-grained system memory.  This allows threadfence_system to work but prevents host memory from being cached on GPU which may have performance impact.");
+
     // Some flags have both compile-time and runtime flags - generate a warning if user enables the runtime flag but the compile-time flag is disabled.
     if (HIP_DB && !COMPILE_HIP_DB) {
         fprintf (stderr, "warning: env var HIP_DB=0x%x but COMPILE_HIP_DB=0.  (perhaps enable COMPILE_HIP_DB in src code before compiling?)\n", HIP_DB);
@@ -1398,6 +1428,7 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream)
 
 void ihipPrintKernelLaunch(const char *kernelName, const grid_launch_parm *lp, const hipStream_t stream)
 {
+
     if (HIP_PROFILE_API || (COMPILE_HIP_DB && HIP_TRACE_API)) {
         std::stringstream os_pre;
         std::stringstream os;
@@ -1515,11 +1546,11 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, g
 //---
 //Called after kernel finishes execution.
 //This releases the lock on the stream.
-void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp)
+void ihipPostLaunchKernel(const char *kernelName, hipStream_t stream, grid_launch_parm &lp)
 {
     tprintf(DB_SYNC, "ihipPostLaunchKernel, unlocking stream\n");
 
-    stream->lockclose_postKernelCommand(lp.av);
+    stream->lockclose_postKernelCommand(kernelName, lp.av);
     MARKER_END();
 }
 
@@ -1883,8 +1914,8 @@ void ihipStream_t::locked_copyAsync(void* dst, const void* src, size_t sizeBytes
             };
 
 
-            if (HIP_LAUNCH_BLOCKING) {
-                tprintf(DB_SYNC, "LAUNCH_BLOCKING for completion of hipMemcpyAsync(%zu)\n", sizeBytes);
+            if (HIP_API_BLOCKING) {
+                tprintf(DB_SYNC, "%s LAUNCH_BLOCKING for completion of hipMemcpyAsync(sz=%zu)\n", ToString(this).c_str(), sizeBytes);
                 this->wait(crit);
             }
 

src/hip_hcc.h

@@ -45,6 +45,7 @@ extern const int release;
 
 // TODO - this blocks both kernels and memory ops.  Perhaps should have separate env var for kernels?
 extern int HIP_LAUNCH_BLOCKING;
+extern int HIP_API_BLOCKING;
 
 extern int HIP_PRINT_ENV;
 extern int HIP_PROFILE_API;
@@ -56,6 +57,8 @@ extern int HIP_STREAM_SIGNALS;  /* number of signals to allocate at stream creat
 extern int HIP_VISIBLE_DEVICES; /* Contains a comma-separated sequence of GPU identifiers */
 extern int HIP_FORCE_P2P_HOST;
 
+extern int HIP_COHERENT_HOST_ALLOC;
+
 
 //---
 // Chicken bits for disabling functionality to work around potential issues:
@@ -156,11 +159,6 @@ extern const char *API_COLOR_END;
 #endif
 
 
-// Compile code that force hipHostMalloc only allocates finegrained system memory.
-#ifndef HIP_COHERENT_HOST_ALLOC
-#define HIP_COHERENT_HOST_ALLOC 0
-#endif
-
 
 
 // Compile support for trace markers that are displayed on CodeXL GUI at start/stop of each function boundary.
@@ -455,7 +453,7 @@ public:
     //---
     // Member functions that begin with locked_ are thread-safe accessors - these acquire / release the critical mutex.
     LockedAccessor_StreamCrit_t  lockopen_preKernelCommand();
-    void                 lockclose_postKernelCommand(hc::accelerator_view *av);
+    void                 lockclose_postKernelCommand(const char *kernelName, hc::accelerator_view *av);
 
 
     void                 locked_wait(bool assertQueueEmpty=false);

src/hip_memory.cpp

@@ -157,20 +157,6 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
     return ihipLogStatus(hip_status);
 }
 
-void ihipReadSingleEnv(int *var_ptr, const char *var_name1, const char *description)
-{
-    char * env = getenv(var_name1);
-
-    // Default is set when variable is initialized (at top of this file), so only override if we find
-    // an environment variable.
-    if (env) {
-        long int v = strtol(env, NULL, 0);
-        *var_ptr = (int) (v);
-    }
-    if (HIP_PRINT_ENV) {
-        printf ("%-30s = %2d : %s\n", var_name1, *var_ptr, description);
-    }
-}
 
 hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
 {
@@ -193,16 +179,12 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
 
         const unsigned supportedFlags = hipHostMallocPortable | hipHostMallocMapped | hipHostMallocWriteCombined;
 
-        // Read from environment variable of HIP_COHERENT_HOST_ALLOC
-        int coherent_alloc=0;
-        ihipReadSingleEnv(&coherent_alloc, "HIP_COHERENT_HOST_ALLOC", "Flag to force allocate finegrained system memory");
-
         if (flags & ~supportedFlags) {
             hip_status = hipErrorInvalidValue;
         }
         else {
             auto device = ctx->getWriteableDevice();
-            if(coherent_alloc){
+            if(HIP_COHERENT_HOST_ALLOC){
                 // Force to allocate finedgrained system memory
                 *ptr = hc::am_alloc(sizeBytes, device->_acc, amHostPinned);
                 if(sizeBytes < 1 && (*ptr == NULL)){
@@ -853,13 +835,13 @@ hipError_t hipMemsetAsync(void* dst, int  value, size_t sizeBytes, hipStream_t s
             }
         }
 
-        stream->lockclose_postKernelCommand(&crit->_av);
+        stream->lockclose_postKernelCommand("hipMemsetAsync", &crit->_av);
 
 
-        if (HIP_LAUNCH_BLOCKING) {
-            tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING wait for memset [stream:%p].\n", __func__, (void*)stream);
+        if (HIP_API_BLOCKING) {
+            tprintf (DB_SYNC, "%s LAUNCH_BLOCKING wait for hipMemsetAsync.\n", ToString(stream).c_str());
             cf.wait();
-            tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING memset completed [stream:%p].\n", __func__, (void*)stream);
+            //tprintf (DB_SYNC, "'%s' LAUNCH_BLOCKING memset completed [stream:%p].\n", __func__, (void*)stream);
         }
     } else {
         e = hipErrorInvalidValue;
@@ -906,7 +888,7 @@ hipError_t hipMemset(void* dst, int  value, size_t sizeBytes )
         // TODO - is hipMemset supposed to be async?
         cf.wait();
 
-        stream->lockclose_postKernelCommand(&crit->_av);
+        stream->lockclose_postKernelCommand("hipMemset", &crit->_av);
 
 
         if (HIP_LAUNCH_BLOCKING) {

src/hip_module.cpp

@@ -333,7 +333,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
 #endif // USE_DISPATCH_HSA_KERNEL
 
 
-        ihipPostLaunchKernel(hStream, lp);
+        ihipPostLaunchKernel(f->_kernelName, hStream, lp);
 
     }