Merge 'master' into 'amd-master'

Change-Id: I1e1c4dd98314f65c5d9a53537eafee324c855718


[ROCm/hip commit: 9fe0bd1a56]

projects/hip/bin/hipcc

@@ -195,7 +195,12 @@ if ($HIP_PLATFORM eq "clang") {
         print ("DEVICE_LIB_PATH=$DEVICE_LIB_PATH\n");
     }
 
-    $HIPCXXFLAGS .= " -std=c++11 -isystem $HIP_CLANG_INCLUDE_PATH";
+    if ($isWindows) {
+      $HIPCXXFLAGS .= " -std=c++14 -fms-extensions -fms-compatibility";
+    } else {
+      $HIPCXXFLAGS .= " -std=c++11";
+    }
+    $HIPCXXFLAGS .= " -isystem $HIP_CLANG_INCLUDE_PATH";
     $HIPLDFLAGS .= " -L$HIP_LIB_PATH";
     if (not $isWindows) {
       $HIPLDFLAGS .= " -Wl,--rpath-link=$HIP_LIB_PATH";
@@ -876,9 +881,7 @@ if ($HIP_PLATFORM eq "clang") {
     }
     $HIP_DEVLIB_FLAGS = " --hip-device-lib-path=$DEVICE_LIB_PATH";
     $HIPCXXFLAGS .= " $HIP_DEVLIB_FLAGS";
-    if ($isWindows) {
+    if (not $isWindows) {
-        $HIPCXXFLAGS .= " -std=c++14 -fms-extensions -fms-compatibility";
-    } else {
         $HIPLDFLAGS .= " -lgcc_s -lgcc -lpthread -lm";
     }
 }

projects/hip/docs/markdown/hip_porting_guide.md

@@ -485,39 +485,7 @@ AMD compilers currently load all data into both the L1 and L2 caches, so __ldg i
 We recommend the following for functional portability:
 
 - For programs that use textures only to benefit from improved caching, use the __ldg instruction
-- Programs that use texture object APIs, work well on HIP
+- Programs that use texture object and reference APIs, work well on HIP
-- For program that use texture reference APIs, use conditional compilation (see [Identify HIP Target Platform](#identify-hip-target-platform)) 
-   - For the `__HIP_PLATFORM_HCC__` path, pass an additional argument to the kernel and in texture fetch API inside kernel as shown below:-
-
-``` 
-texture<float, 2, hipReadModeElementType> tex;
-
-__global__ void tex2DKernel(float* outputData,
-#ifdef __HIP_PLATFORM_HCC__
-                             hipTextureObject_t textureObject,
-#endif
-                             int width,
-                             int height)
-{
-    int x = blockIdx.x*blockDim.x + threadIdx.x;
-    int y = blockIdx.y*blockDim.y + threadIdx.y;
-#ifdef __HIP_PLATFORM_HCC__
-    outputData[y*width + x] = tex2D(tex, textureObject, x, y);
-#else
-    outputData[y*width + x] = tex2D(tex, x, y);
-#endif
-}
-
-// Host code:
-void myFunc () 
-{
-    // ...
-
-#ifdef __HIP_PLATFORM_HCC__
-    hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, tex.textureObject, width, height);
-#else
-    hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, width, height);
-#endif
 
 
 ``` 

projects/hip/hipify-clang/src/CUDA2HIP_Runtime_API_types.cpp

@@ -1067,7 +1067,7 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
   // CU_SHAREDMEM_CARVEOUT_MAX_SHARED
   {"cudaSharedmemCarveoutMaxShared",                                   {"hipSharedmemCarveoutMaxShared",                            "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 100
   // CU_SHAREDMEM_CARVEOUT_MAX_L1
-  {"cudaSharedmemCarveoutMaxShared",                                   {"hipSharedmemCarveoutMaxL1",                                "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 0
+  {"cudaSharedmemCarveoutMaxL1",                                       {"hipSharedmemCarveoutMaxL1",                                "", CONV_NUMERIC_LITERAL, API_RUNTIME, HIP_UNSUPPORTED}}, // 0
 
   // CUsharedconfig
   {"cudaSharedMemConfig",                                              {"hipSharedMemConfig",                                       "", CONV_TYPE, API_RUNTIME}},
@@ -1244,6 +1244,9 @@ const std::map<llvm::StringRef, hipCounter> CUDA_RUNTIME_TYPE_NAME_MAP {
   // CUtexObject
   {"cudaTextureObject_t",                                              {"hipTextureObject_t",                                       "", CONV_TYPE, API_RUNTIME}},
 
+  // CUuuid
+  {"cudaUUID_t",                                                       {"hipUUID_t",                                                "", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}},
+
   // 5. Defines
 
   // no analogue

projects/hip/include/hip/hcc_detail/functional_grid_launch.hpp

@@ -151,6 +151,20 @@ void hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* blockSize,
                                       dynSharedMemPerBlk, blockSizeLimit);
 }
 
+template <typename F>
+inline
+void hipOccupancyMaxActiveBlocksPerMultiprocessor(uint32_t* numBlocks, F kernel,
+    uint32_t blockSize, size_t dynSharedMemPerBlk) {
+
+    using namespace hip_impl;
+
+    hip_impl::hip_init();
+    auto f = get_program_state().kernel_descriptor(reinterpret_cast<std::uintptr_t>(kernel),
+                                                   target_agent(0));
+
+    hipOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, f, blockSize, dynSharedMemPerBlk);
+}
+
 template <typename... Args, typename F = void (*)(Args...)>
 inline
 void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,

projects/hip/include/hip/hcc_detail/hip_runtime_api.h

@@ -2749,7 +2749,7 @@ hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsLi
  *
  * @param [out] gridSize           minimum grid size for maximum potential occupancy
  * @param [out] blockSize          block size for maximum potential occupancy
- * @param [in]  f                  kernel to launch
+ * @param [in]  f                  kernel function for which occupancy is calulated
  * @param [in]  dynSharedMemPerBlk dynamic shared memory usage (in bytes) intended for each block
  * @param [in]  blockSizeLimit     the maximum block size for the kernel, use 0 for no limit
  *
@@ -2765,10 +2765,10 @@ hipError_t hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* block
  * @param [out] numBlocks        Returned occupancy
  * @param [in]  func             Kernel function for which occupancy is calulated
  * @param [in]  blockSize        Block size the kernel is intended to be launched with
- * @param [in]  dynamicSMemSize  Per - block dynamic shared memory usage intended, in bytes
+ * @param [in]  dynSharedMemPerBlk dynamic shared memory usage (in bytes) intended for each block
  */
 hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
-   int* numBlocks, const void* f, int  blockSize, size_t dynamicSMemSize);
+   uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk);
 
 /**
  * @brief Returns occupancy for a device function.
@@ -2776,11 +2776,11 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
  * @param [out] numBlocks        Returned occupancy
  * @param [in]  func             Kernel function for which occupancy is calulated
  * @param [in]  blockSize        Block size the kernel is intended to be launched with
- * @param [in]  dynamicSMemSize  Per - block dynamic shared memory usage intended, in bytes
+ * @param [in]  dynSharedMemPerBlk dynamic shared memory usage (in bytes) intended for each block
  * @param [in]  flags            Extra flags for occupancy calculation (currently ignored)
  */
 hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
-   int* numBlocks, const void* f, int  blockSize, size_t dynamicSMemSize, unsigned int flags);
+   uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk, unsigned int flags);
 
 /**
  * @brief Launches kernels on multiple devices and guarantees all specified kernels are dispatched
@@ -3123,19 +3123,6 @@ hipError_t hipBindTextureToMipmappedArray(const texture<T, dim, readMode>& tex,
     return hipSuccess;
 }
 
-template <class T>
-inline hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
-    int* numBlocks, T f, int  blockSize, size_t dynamicSMemSize) {
-    return hipOccupancyMaxActiveBlocksPerMultiprocessor(
-        numBlocks, reinterpret_cast<const void*>(f), blockSize, dynamicSMemSize);
-}
-
-template <class T>
-inline hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
-    int* numBlocks, T f, int  blockSize, size_t dynamicSMemSize, unsigned int flags) {
-    return hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
-        numBlocks, reinterpret_cast<const void*>(f), blockSize, dynamicSMemSize, flags);
-}
 
 template <class T>
 inline hipError_t hipLaunchCooperativeKernel(T f, dim3 gridDim, dim3 blockDim,

projects/hip/include/hip/hip_runtime_api.h

@@ -117,10 +117,12 @@ typedef struct hipDeviceProp_t {
     int integrated;            ///< APU vs dGPU
     int cooperativeLaunch;            ///< HIP device supports cooperative launch
     int cooperativeMultiDeviceLaunch; ///< HIP device supports cooperative launch on multiple devices
-#if !__HIP_VDI__ // Temporarily disable the following three new fields for HIP/VDI runtime
+#if !__HIP_VDI__ // Temporarily disable the following five new fields for HIP/VDI runtime
     int maxTexture1D;          ///< Maximum number of elements in 1D images
     int maxTexture2D[2];       ///< Maximum dimensions (width, height) of 2D images, in image elements
     int maxTexture3D[3];       ///< Maximum dimensions (width, height, depth) of 3D images, in image elements
+    unsigned int* hdpMemFlushCntl;      ///< Addres of HDP_MEM_COHERENCY_FLUSH_CNTL register
+    unsigned int* hdpRegFlushCntl;      ///< Addres of HDP_REG_COHERENCY_FLUSH_CNTL register
 #endif
 } hipDeviceProp_t;
 
@@ -306,8 +308,10 @@ typedef enum hipDeviceAttribute_t {
     hipDeviceAttributeMaxTexture2DHeight,   ///< Maximum dimension height of 2D images in image elements
     hipDeviceAttributeMaxTexture3DWidth,    ///< Maximum dimension width of 3D images in image elements
     hipDeviceAttributeMaxTexture3DHeight,   ///< Maximum dimensions height of 3D images in image elements
-    hipDeviceAttributeMaxTexture3DDepth     ///< Maximum dimensions depth of 3D images in image elements
+    hipDeviceAttributeMaxTexture3DDepth,    ///< Maximum dimensions depth of 3D images in image elements
 
+    hipDeviceAttributeHdpMemFlushCntl,      ///< Address of the HDP_MEM_COHERENCY_FLUSH_CNTL register
+    hipDeviceAttributeHdpRegFlushCntl       ///< Address of the HDP_REG_COHERENCY_FLUSH_CNTL register
 } hipDeviceAttribute_t;
 
 enum hipComputeMode {

projects/hip/include/hip/nvcc_detail/hip_runtime_api.h

@@ -715,17 +715,20 @@ inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t* p_prop, int dev
     struct cudaDeviceProp cdprop;
     cudaError_t cerror;
     cerror = cudaGetDeviceProperties(&cdprop, device);
+
     strncpy(p_prop->name, cdprop.name, 256);
     p_prop->totalGlobalMem = cdprop.totalGlobalMem;
     p_prop->sharedMemPerBlock = cdprop.sharedMemPerBlock;
     p_prop->regsPerBlock = cdprop.regsPerBlock;
     p_prop->warpSize = cdprop.warpSize;
+    p_prop->maxThreadsPerBlock = cdprop.maxThreadsPerBlock;
     for (int i = 0; i < 3; i++) {
         p_prop->maxThreadsDim[i] = cdprop.maxThreadsDim[i];
         p_prop->maxGridSize[i] = cdprop.maxGridSize[i];
     }
-    p_prop->maxThreadsPerBlock = cdprop.maxThreadsPerBlock;
     p_prop->clockRate = cdprop.clockRate;
+    p_prop->memoryClockRate = cdprop.memoryClockRate;
+    p_prop->memoryBusWidth = cdprop.memoryBusWidth;
     p_prop->totalConstMem = cdprop.totalConstMem;
     p_prop->major = cdprop.major;
     p_prop->minor = cdprop.minor;
@@ -733,44 +736,38 @@ inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t* p_prop, int dev
     p_prop->l2CacheSize = cdprop.l2CacheSize;
     p_prop->maxThreadsPerMultiProcessor = cdprop.maxThreadsPerMultiProcessor;
     p_prop->computeMode = cdprop.computeMode;
-    p_prop->canMapHostMemory = cdprop.canMapHostMemory;
+    p_prop->clockInstructionRate = cdprop.clockRate; // Same as clock-rate:
-    p_prop->memoryClockRate = cdprop.memoryClockRate;
-    p_prop->memoryBusWidth = cdprop.memoryBusWidth;
-
-    // Same as clock-rate:
-    p_prop->clockInstructionRate = cdprop.clockRate;
 
     int ccVers = p_prop->major * 100 + p_prop->minor * 10;
-
     p_prop->arch.hasGlobalInt32Atomics = (ccVers >= 110);
     p_prop->arch.hasGlobalFloatAtomicExch = (ccVers >= 110);
     p_prop->arch.hasSharedInt32Atomics = (ccVers >= 120);
     p_prop->arch.hasSharedFloatAtomicExch = (ccVers >= 120);
-
     p_prop->arch.hasFloatAtomicAdd = (ccVers >= 200);
-
     p_prop->arch.hasGlobalInt64Atomics = (ccVers >= 120);
     p_prop->arch.hasSharedInt64Atomics = (ccVers >= 110);
-
     p_prop->arch.hasDoubles = (ccVers >= 130);
-
     p_prop->arch.hasWarpVote = (ccVers >= 120);
     p_prop->arch.hasWarpBallot = (ccVers >= 200);
     p_prop->arch.hasWarpShuffle = (ccVers >= 300);
     p_prop->arch.hasFunnelShift = (ccVers >= 350);
-
     p_prop->arch.hasThreadFenceSystem = (ccVers >= 200);
     p_prop->arch.hasSyncThreadsExt = (ccVers >= 200);
-
     p_prop->arch.hasSurfaceFuncs = (ccVers >= 200);
     p_prop->arch.has3dGrid = (ccVers >= 200);
     p_prop->arch.hasDynamicParallelism = (ccVers >= 350);
 
     p_prop->concurrentKernels = cdprop.concurrentKernels;
+    p_prop->pciDomainID = cdprop.pciDomainID;
     p_prop->pciBusID = cdprop.pciBusID;
     p_prop->pciDeviceID = cdprop.pciDeviceID;
     p_prop->maxSharedMemoryPerMultiProcessor = cdprop.sharedMemPerMultiprocessor;
     p_prop->isMultiGpuBoard = cdprop.isMultiGpuBoard;
+    p_prop->canMapHostMemory = cdprop.canMapHostMemory;
+    p_prop->gcnArch = 0; // Not a GCN arch
+    p_prop->integrated = cdprop.integrated;
+    p_prop->cooperativeLaunch = cdprop.cooperativeLaunch;
+    p_prop->cooperativeMultiDeviceLaunch = cdprop.cooperativeMultiDeviceLaunch;
 
     p_prop->maxTexture1D    = cdprop.maxTexture1D;
     p_prop->maxTexture2D[0] = cdprop.maxTexture2D[0];

projects/hip/src/hip_device.cpp

@@ -293,6 +293,17 @@ hipError_t ihipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device
                 break;
             case hipDeviceAttributeMaxTexture3DDepth:
                 *pi = prop->maxTexture3D[2];
+            case hipDeviceAttributeHdpMemFlushCntl:
+                {
+                    uint32_t** hdp = reinterpret_cast<uint32_t**>(pi);
+                    *hdp = prop->hdpMemFlushCntl;
+                }
+                break;
+            case hipDeviceAttributeHdpRegFlushCntl:
+                {
+                    uint32_t** hdp = reinterpret_cast<uint32_t**>(pi);
+                    *hdp = prop->hdpRegFlushCntl;
+                }
                 break;
             default:
                 e = hipErrorInvalidValue;

projects/hip/src/hip_event.cpp

@@ -191,10 +191,8 @@ hipError_t hipEventSynchronize(hipEvent_t event) {
             ctx->locked_syncDefaultStream(true, true);
             return ihipLogStatus(hipSuccess);
         } else {
-            ecd._stream->locked_eventWaitComplete(
+            ecd.marker().wait((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked
-                ecd.marker(), (event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked
                                                                      : hc::hcWaitModeActive);
-
             return ihipLogStatus(hipSuccess);
         }
     } else {

projects/hip/src/hip_hcc.cpp

@@ -316,16 +316,20 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t& crit) {
     tprintf(DB_SYNC, "%s wait for queue-empty..\n", ToString(this).c_str());
 
     crit->_av.wait(waitMode());
-
-    crit->_kernelCnt = 0;
 }
 
 //---
 // Wait for all kernel and data copy commands in this stream to complete.
 void ihipStream_t::locked_wait() {
-    LockedAccessor_StreamCrit_t crit(_criticalData);
+    // create a marker while holding stream lock,
+    // but release lock prior to waiting on the marker
+    hc::completion_future marker;
+    {
+        LockedAccessor_StreamCrit_t crit(_criticalData);
+        marker = crit->_av.create_marker(hc::no_scope);
+    }
 
-    wait(crit);
+    marker.wait(waitMode());
 };
 
 // Causes current stream to wait for specified event to complete:
@@ -340,30 +344,14 @@ void ihipStream_t::locked_streamWaitEvent(ihipEventData_t& ecd) {
 // Causes current stream to wait for specified event to complete:
 // Note this does not provide any kind of host serialization.
 bool ihipStream_t::locked_eventIsReady(hipEvent_t event) {
-    // Event query that returns "Complete" may cause HCC to manipulate
-    // internal queue state so lock the stream's queue here.
-    LockedAccessor_StreamCrit_t scrit(_criticalData);
-
     LockedAccessor_EventCrit_t ecrit(event->criticalData());
 
     return (ecrit->_eventData.marker().is_ready());
 }
 
-// Waiting on event can cause HCC to reclaim stream resources - so need to lock the stream.
-void ihipStream_t::locked_eventWaitComplete(hc::completion_future& marker,
-                                            hc::hcWaitMode waitMode) {
-    LockedAccessor_StreamCrit_t crit(_criticalData);
-
-    marker.wait(waitMode);
-}
-
-
 // Create a marker in this stream.
 // Save state in the event so it can track the status of the event.
 hc::completion_future ihipStream_t::locked_recordEvent(hipEvent_t event) {
-    // Lock the stream to prevent simultaneous access
-    LockedAccessor_StreamCrit_t crit(_criticalData);
-
     auto scopeFlag = hc::accelerator_scope;
     // The env var HIP_EVENT_SYS_RELEASE sets the default,
     // The explicit flags override the env var (if specified)
@@ -375,6 +363,8 @@ hc::completion_future ihipStream_t::locked_recordEvent(hipEvent_t event) {
         scopeFlag = HIP_EVENT_SYS_RELEASE ? hc::system_scope : hc::accelerator_scope;
     }
 
+    // Lock the stream to prevent simultaneous access
+    LockedAccessor_StreamCrit_t crit(_criticalData);
     return crit->_av.create_marker(scopeFlag);
 };
 
@@ -925,6 +915,15 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) {
     err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_EXT_AGENT_INFO_IMAGE_3D_MAX_ELEMENTS,
           prop->maxTexture3D);
     DeviceErrorCheck(err);
+
+    // Get Agent HDP Flush Register Memory
+    hsa_amd_hdp_flush_t hdpinfo;
+    err = hsa_agent_get_info(_hsaAgent, (hsa_agent_info_t)HSA_AMD_AGENT_INFO_HDP_FLUSH, &hdpinfo);
+    DeviceErrorCheck(err);
+
+    prop->hdpMemFlushCntl = hdpinfo.HDP_MEM_FLUSH_CNTL;
+    prop->hdpRegFlushCntl = hdpinfo.HDP_REG_FLUSH_CNTL;
+
     return e;
 }
 

projects/hip/src/hip_hcc_internal.h

@@ -476,7 +476,7 @@ template <typename MUTEX_TYPE>
 class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE> {
    public:
     ihipStreamCriticalBase_t(ihipStream_t* parentStream, hc::accelerator_view av)
-        : _kernelCnt(0), _av(av), _parent(parentStream){};
+        : _av(av), _parent(parentStream){};
 
     ~ihipStreamCriticalBase_t() {}
 
@@ -500,7 +500,6 @@ class ihipStreamCriticalBase_t : public LockedBase<MUTEX_TYPE> {
 
    public:
     ihipStream_t* _parent;
-    uint32_t _kernelCnt;  // Count of inflight kernels in this stream.  Reset at ::wait().
 
     hc::accelerator_view _av;
 
@@ -564,7 +563,6 @@ class ihipStream_t {
     hc::completion_future locked_recordEvent(hipEvent_t event);
 
     bool locked_eventIsReady(hipEvent_t event);
-    void locked_eventWaitComplete(hc::completion_future& marker, hc::hcWaitMode waitMode);
 
     ihipStreamCritical_t& criticalData() { return _criticalData; };
 

projects/hip/src/hip_memory.cpp

@@ -198,7 +198,7 @@ hipError_t hipPointerGetAttributes(hipPointerAttribute_t* attributes, const void
             attributes->isManaged = 0;
             attributes->allocationFlags = 0;
 
-            e = hipErrorUnknown;  // TODO - should be hipErrorInvalidValue ?
+            e = hipErrorInvalidValue;
         }
     }
     return ihipLogStatus(e);

projects/hip/src/hip_module.cpp

@@ -318,16 +318,16 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
 
 hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
                                               int  numDevices, unsigned int  flags) {
-
+    HIP_INIT_API(hipExtLaunchMultiKernelMultiDevice, launchParamsList,numDevices,flags);
     hipError_t result;
 
     if ((numDevices > g_deviceCnt) || (launchParamsList == nullptr)) {
-        return hipErrorInvalidValue;
+        return ihipLogStatus(hipErrorInvalidValue);
     }
 
     hipFunction_t* kds = reinterpret_cast<hipFunction_t*>(malloc(sizeof(hipFunction_t) * numDevices));
     if (kds == nullptr) {
-        return hipErrorNotInitialized;
+        return ihipLogStatus(hipErrorNotInitialized);
     }
 
     // prepare all kernel descriptors for each device as all streams will be locked in the next loop
@@ -335,13 +335,13 @@ hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
         const hipLaunchParams& lp = launchParamsList[i];
         if (lp.stream == nullptr) {
             free(kds);
-            return hipErrorNotInitialized;
+            return ihipLogStatus(hipErrorNotInitialized);
         }
         kds[i] = hip_impl::get_program_state().kernel_descriptor(reinterpret_cast<std::uintptr_t>(lp.func),
                 hip_impl::target_agent(lp.stream));
         if (kds[i] == nullptr) {
             free(kds);
-            return hipErrorInvalidValue;
+            return ihipLogStatus(hipErrorInvalidValue);
         }
         hip_impl::kernargs_size_align kargs = hip_impl::get_program_state().get_kernargs_size_align(
                 reinterpret_cast<std::uintptr_t>(lp.func));
@@ -382,7 +382,7 @@ hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList,
 
     free(kds);
 
-    return result;
+    return ihipLogStatus(result);
 }
 
 namespace hip_impl {
@@ -785,6 +785,7 @@ hipFuncAttributes make_function_attributes(const ihipModuleSymbol_t& kd) {
 
 hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func)
 {
+    HIP_INIT_API(hipFuncGetAttributes, attr, func);
     using namespace hip_impl;
 
     if (!attr) return hipErrorInvalidValue;
@@ -797,7 +798,7 @@ hipError_t hipFuncGetAttributes(hipFuncAttributes* attr, const void* func)
 
     *attr = make_function_attributes(*kd);
 
-    return hipSuccess;
+    return ihipLogStatus(hipSuccess);
 }
 
 hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) {
@@ -878,6 +879,30 @@ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const
     return ihipLogStatus(hipSuccess);
 }
 
+void getGprsLdsUsage(hipFunction_t f, size_t* usedVGPRS, size_t* usedSGPRS, size_t* usedLDS)
+{
+    bool is_code_object_v3 = f->_name.find(".kd") != std::string::npos;
+    if (is_code_object_v3) {
+        const auto header = reinterpret_cast<const amd_kernel_code_v3_t*>(f->_header);
+        // GRANULATED_WAVEFRONT_VGPR_COUNT is specified in 0:5 bits of COMPUTE_PGM_RSRC1
+        // the granularity for gfx6-gfx9 is max(0, ceil(vgprs_used / 4) - 1)
+        *usedVGPRS = ((header->compute_pgm_rsrc1 & 0x3F) + 1) << 2;
+        // GRANULATED_WAVEFRONT_SGPR_COUNT is specified in 6:9 bits of COMPUTE_PGM_RSRC1
+        // the granularity for gfx9+ is 2 * max(0, ceil(sgprs_used / 16) - 1)
+        *usedSGPRS = ((((header->compute_pgm_rsrc1 & 0x3C0) >> 6) >> 1) + 1) << 4;
+        *usedLDS = header->group_segment_fixed_size;
+    }
+    else {
+        const auto header = f->_header;
+        // VGPRs granularity is 4
+        *usedVGPRS = ((header->workitem_vgpr_count + 3) >> 2) << 2;
+        // adding 2 to take into account the 2 VCC registers & handle the granularity of 16
+        *usedSGPRS = header->wavefront_sgpr_count + 2;
+        *usedSGPRS = ((*usedSGPRS + 15) >> 4) << 4;
+        *usedLDS = header->workgroup_group_segment_byte_size;
+    }
+}
+
 hipError_t ihipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* blockSize,
                                               hipFunction_t f, size_t dynSharedMemPerBlk,
                                               uint32_t blockSizeLimit)
@@ -885,10 +910,8 @@ hipError_t ihipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* bloc
     using namespace hip_impl;
 
     auto ctx = ihipGetTlsDefaultCtx();
-    hipError_t ret = hipSuccess;
-
     if (ctx == nullptr) {
-        ret = hipErrorInvalidDevice;
+        return hipErrorInvalidDevice;
     }
 
     hipDeviceProp_t prop{};
@@ -899,26 +922,7 @@ hipError_t ihipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* bloc
     size_t usedVGPRS = 0;
     size_t usedSGPRS = 0;
     size_t usedLDS = 0;
-    bool is_code_object_v3 = f->_name.find(".kd") != std::string::npos;
+    getGprsLdsUsage(f, &usedVGPRS, &usedSGPRS, &usedLDS);
-    if (is_code_object_v3) {
-        const auto header = reinterpret_cast<const amd_kernel_code_v3_t*>(f->_header);
-        // GRANULATED_WAVEFRONT_VGPR_COUNT is specified in 0:5 bits of COMPUTE_PGM_RSRC1
-        // the granularity for gfx6-gfx9 is max(0, ceil(vgprs_used / 4) - 1)
-        usedVGPRS = ((header->compute_pgm_rsrc1 & 0x3F) + 1) << 2;
-        // GRANULATED_WAVEFRONT_SGPR_COUNT is specified in 6:9 bits of COMPUTE_PGM_RSRC1
-        // the granularity for gfx9+ is 2 * max(0, ceil(sgprs_used / 16) - 1)
-        usedSGPRS = ((((header->compute_pgm_rsrc1 & 0x3C0) >> 6) >> 1) + 1) << 4;
-        usedLDS = header->group_segment_fixed_size;
-    }
-    else {
-        const auto header = f->_header;
-        // VGPRs granularity is 4
-        usedVGPRS = ((header->workitem_vgpr_count + 3) >> 2) << 2;
-        // adding 2 to take into account the 2 VCC registers & handle the granularity of 16
-        usedSGPRS = header->wavefront_sgpr_count + 2;
-        usedSGPRS = ((usedSGPRS + 15) >> 4) << 4;
-        usedLDS = header->workgroup_group_segment_byte_size;
-    }
 
     // try different workgroup sizes to find the maximum potential occupancy
     // based on the usage of VGPRs and LDS
@@ -1008,10 +1012,9 @@ hipError_t ihipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* bloc
     *blockSize = maxWavefronts * wavefrontSize;
     *gridSize = min((maxThreadsCnt + *blockSize - 1) / *blockSize, prop.multiProcessorCount);
 
-    return ret;
+    return hipSuccess;
 }
 
-
 hipError_t hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* blockSize,
                                              hipFunction_t f, size_t dynSharedMemPerBlk,
                                              uint32_t blockSizeLimit)
@@ -1021,3 +1024,74 @@ hipError_t hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* block
     return ihipLogStatus(ihipOccupancyMaxPotentialBlockSize(
         gridSize, blockSize, f, dynSharedMemPerBlk, blockSizeLimit));
 }
+
+hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor(
+   uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk)
+{
+    using namespace hip_impl;
+
+    auto ctx = ihipGetTlsDefaultCtx();
+    if (ctx == nullptr) {
+        return hipErrorInvalidDevice;
+    }
+
+    hipDeviceProp_t prop{};
+    ihipGetDeviceProperties(&prop, ihipGetTlsDefaultCtx()->getDevice()->_deviceId);
+
+    prop.regsPerBlock = prop.regsPerBlock ? prop.regsPerBlock : 64 * 1024;
+
+    size_t usedVGPRS = 0;
+    size_t usedSGPRS = 0;
+    size_t usedLDS = 0;
+    getGprsLdsUsage(f, &usedVGPRS, &usedSGPRS, &usedLDS);
+
+    // Due to SPI and private memory limitations, the max of wavefronts per CU in 32
+    size_t wavefrontSize = prop.warpSize;
+    size_t maxWavefrontsPerCU = min(prop.maxThreadsPerMultiProcessor / wavefrontSize, 32);
+
+    const size_t simdPerCU = 4;
+    const size_t maxWavesPerSimd = maxWavefrontsPerCU / simdPerCU;
+
+    size_t numWavefronts = (blockSize + wavefrontSize - 1) / wavefrontSize;
+
+    size_t availableVGPRs = (prop.regsPerBlock / wavefrontSize / simdPerCU);
+    size_t vgprs_alu_occupancy = simdPerCU * std::min(maxWavesPerSimd, availableVGPRs / usedVGPRS);
+
+    // Calculate blocks occupancy per CU based on VGPR usage
+    *numBlocks = vgprs_alu_occupancy / numWavefronts;
+
+    const size_t availableSGPRs = (prop.gcnArch < 800) ? 512 : 800;
+    size_t sgprs_alu_occupancy = simdPerCU * ((usedSGPRS == 0) ? maxWavesPerSimd
+        : std::min(maxWavesPerSimd, availableSGPRs / usedSGPRS));
+
+    // Calculate blocks occupancy per CU based on SGPR usage
+    *numBlocks = std::min(*numBlocks, (uint32_t) (sgprs_alu_occupancy / numWavefronts)); 
+
+    size_t total_used_lds = usedLDS + dynSharedMemPerBlk;
+    if (total_used_lds != 0) {
+      // Calculate LDS occupacy per CU. lds_per_cu / (static_lsd + dynamic_lds)
+      size_t lds_occupancy = prop.maxSharedMemoryPerMultiProcessor / total_used_lds;
+      *numBlocks = std::min(*numBlocks, (uint32_t) lds_occupancy);
+    }
+
+    return hipSuccess;
+}
+
+hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
+   uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk)
+{
+    HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessor, numBlocks, f, blockSize, dynSharedMemPerBlk);
+
+    return ihipLogStatus(ihipOccupancyMaxActiveBlocksPerMultiprocessor(
+        numBlocks, f, blockSize, dynSharedMemPerBlk));
+}
+
+hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
+   uint32_t* numBlocks, hipFunction_t f, uint32_t  blockSize, size_t dynSharedMemPerBlk,
+   unsigned int flags)
+{
+    HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags, numBlocks, f, blockSize, dynSharedMemPerBlk, flags);
+
+    return ihipLogStatus(ihipOccupancyMaxActiveBlocksPerMultiprocessor(
+        numBlocks, f, blockSize, dynSharedMemPerBlk));
+}

projects/hip/src/hip_stream.cpp

@@ -142,9 +142,8 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int
                 // conservative wait on host for the specified event to complete:
                 // return _stream->locked_eventWaitComplete(this, waitMode);
                 //
-                ecd._stream->locked_eventWaitComplete(
+                ecd.marker().wait((event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked
-                ecd.marker(), (event->_flags & hipEventBlockingSync) ? hc::hcWaitModeBlocked
+                                                                         : hc::hcWaitModeActive);
-                                                                     : hc::hcWaitModeActive);
             } else {
                 stream = ihipSyncAndResolveStream(stream);
                 // This will use create_blocking_marker to wait on the specified queue.

projects/hip/tests/src/deviceLib/hip_mbcnt.cpp

@@ -59,12 +59,12 @@ int main() {
 
     cout << "hip Device prop succeeded " << endl;
 
-    constexpr unsigned int wave_size = 64;
     constexpr unsigned int num_waves_per_block = 2;
-    constexpr unsigned int num_threads_per_block = wave_size * num_waves_per_block;
+    const unsigned int wave_size = devProp.warpSize;
-    constexpr unsigned int num_blocks = 2;
+    const unsigned int num_threads_per_block = wave_size * num_waves_per_block;
-    constexpr unsigned int num_threads = num_threads_per_block * num_blocks;
+    const unsigned int num_blocks = 2;
-    constexpr size_t buffer_size = num_threads * sizeof(unsigned int);
+    const unsigned int num_threads = num_threads_per_block * num_blocks;
+    const size_t buffer_size = num_threads * sizeof(unsigned int);
 
     HIP_ASSERT(hipMalloc((void**)&device_mbcnt_lo, buffer_size));
     HIP_ASSERT(hipMalloc((void**)&device_mbcnt_hi, buffer_size));

projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp

@@ -56,6 +56,22 @@ hipError_t test_hipDeviceGetAttribute(int deviceId, hipDeviceAttribute_t attr,
     return hipSuccess;
 }
 
+hipError_t test_hipDeviceGetHdpAddress(int deviceId, hipDeviceAttribute_t attr,
+                                       uint32_t* expectedValue = (uint32_t*)0xdeadbeef) {
+    uint32_t* value = 0;
+    std::cout << "Test hipDeviceGetHdpAddress attribute " << attr;
+    if (expectedValue != (uint32_t*)0xdeadbeef) {
+        std::cout << " expected value " << expectedValue;
+    }
+    hipError_t e = hipDeviceGetAttribute((int*) &value, attr, deviceId);
+    std::cout << " actual value " << value << std::endl;
+    if ((expectedValue != (uint32_t*)0xdeadbeef) && value != expectedValue) {
+        std::cout << "fail" << std::endl;
+        return hipErrorInvalidValue;
+    }
+    return hipSuccess;
+}
+
 int main(int argc, char* argv[]) {
     int deviceId;
     CHECK(hipGetDevice(&deviceId));
@@ -116,5 +132,9 @@ int main(int argc, char* argv[]) {
     CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture3DWidth, props.maxTexture3D[0]));
     CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture3DHeight, props.maxTexture3D[1]));
     CHECK(test_hipDeviceGetAttribute(deviceId, hipDeviceAttributeMaxTexture3DDepth, props.maxTexture3D[2]));
+#ifndef __HIP_PLATFORM_NVCC__
+    CHECK(test_hipDeviceGetHdpAddress(deviceId, hipDeviceAttributeHdpMemFlushCntl, props.hdpMemFlushCntl));
+    CHECK(test_hipDeviceGetHdpAddress(deviceId, hipDeviceAttributeHdpRegFlushCntl, props.hdpRegFlushCntl));
+#endif
     passed();
 };

projects/hip/tests/src/runtimeApi/module/hipFuncGetAttributes.cpp

@@ -39,7 +39,6 @@ void fn(float* px, float* py)
 }
 
 int main() {
-    hipInit(0);
 
     hipFuncAttributes attr{};
 

projects/hip/tests/src/runtimeApi/module/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp

@@ -0,0 +1,78 @@
+/*
+Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+// Test the Grid_Launch syntax.
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t
+ * HIT_END
+ */
+
+#include "hip/hip_runtime.h"
+#include "test_common.h"
+
+#define fileName "vcpy_kernel.code"
+#define kernel_name "hello_world"
+
+
+__global__ void f1(float *a) { *a = 1.0; }
+
+template <typename T>
+__global__ void f2(T *a) { *a = 1; }
+
+
+
+int main(int argc, char* argv[]) {
+
+    // test case for using kernel function pointer
+    uint32_t gridSize = 0;
+    uint32_t blockSize = 0;
+    hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0);
+    assert(gridSize != 0 && blockSize != 0);
+
+    uint32_t numBlock = 0;
+    hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, blockSize, 0);
+    assert(numBlock != 0);
+
+
+    // test case for using kernel function pointer with template
+    gridSize = 0;
+    blockSize = 0;
+    hipOccupancyMaxPotentialBlockSize<void(*)(int *)>(&gridSize, &blockSize, f2, 0, 0);
+    assert(gridSize != 0 && blockSize != 0);
+
+    numBlock = 0;
+    hipOccupancyMaxActiveBlocksPerMultiprocessor<void(*)(int *)>(&numBlock, f2, blockSize, 0);
+    assert(numBlock != 0);
+
+
+    // test case for using kernel with hipFunction_t type
+    numBlock = 0;
+    hipModule_t Module;
+    hipFunction_t Function;
+    HIPCHECK(hipModuleLoad(&Module, fileName));
+    HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
+    HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, Function, blockSize, 0));
+    assert(numBlock != 0);
+
+    passed();
+}