diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 55c7cfd847..1a40f6bd21 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -95,10 +95,12 @@ enum hipLimit_t #define hipHostRegisterIoMemory 0x4 ///< Not supported. -#define hipDeviceScheduleAuto 0x0 -#define hipDeviceScheduleSpin 0x1 -#define hipDeviceScheduleYield 0x2 -#define hipDeviceBlockingSync 0x4 +#define hipDeviceScheduleAuto 0x0 ///< Automatically select between Spin and Yield +#define hipDeviceScheduleSpin 0x1 ///< Dedicate a CPU core to spin-wait. Provides lowest latency, but burns a CPU core and may consume more power. +#define hipDeviceScheduleYield 0x2 ///< Yield the CPU to the operating system when waiting. May increase latency, but lowers power and is friendlier to other threads in the system. +#define hipDeviceScheduleBlockingSync 0x4 +#define hipDeviceScheduleMask 0x7 + #define hipDeviceMapHost 0x8 #define hipDeviceLmemResizeToMax 0x16 @@ -383,9 +385,18 @@ hipError_t hipDeviceSetSharedMemConfig ( hipSharedMemConfig config ); * * @param [in] flags * + * The schedule flags impact how HIP waits for the completion of a command running on a device. + * hipDeviceScheduleSpin : HIP runtime will actively spin in the thread which submitted the work until the command completes. This offers the lowest latency, but will consume a CPU core and may increase power. + * hipDeviceScheduleYield : The HIP runtime will yield the CPU to system so that other tasks can use it. This may increase latency to detect the completion but will consume less power and is friendlier to other tasks in the system. + * hipDeviceScheduleBlockingSync : On ROCm platform, this is a synonym for hipDeviceScheduleYield. + * hipDeviceScheduleAuto : Use a hueristic to select between Spin and Yield modes. If the number of HIP contexts is greater than the number of logical processors in the system, use Spin scheduling. Else use Yield scheduling. + * + * + * hipDeviceMapHost : Allow mapping host memory. On ROCM, this is always allowed and the flag is ignored. + * hipDeviceLmemResizeToMax : @warning ROCm silently ignores this flag. + * * @returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorSetOnActiveProcess * - * Note: Only hipDeviceScheduleAuto and hipDeviceMapHost are supported * */ hipError_t hipSetDeviceFlags ( unsigned flags); diff --git a/projects/clr/hipamd/src/hip_device.cpp b/projects/clr/hipamd/src/hip_device.cpp index fda5724d7f..7168ccfbe6 100644 --- a/projects/clr/hipamd/src/hip_device.cpp +++ b/projects/clr/hipamd/src/hip_device.cpp @@ -268,7 +268,7 @@ hipError_t hipSetDeviceFlags( unsigned int flags) { HIP_INIT_API(flags); - hipError_t e; + hipError_t e = hipSuccess; auto * ctx = ihipGetTlsDefaultCtx(); @@ -276,7 +276,25 @@ hipError_t hipSetDeviceFlags( unsigned int flags) // TODO : Review error handling behavior for this function, it often returns ErrorSetOnActiveProcess if (ctx) { ctx->_ctxFlags = ctx->_ctxFlags | flags; - e = hipSuccess; + if (flags & hipDeviceScheduleMask) { + switch (hipDeviceScheduleMask) { + case hipDeviceScheduleAuto: + case hipDeviceScheduleSpin: + case hipDeviceScheduleYield: + case hipDeviceScheduleBlockingSync: + e = hipSuccess; + break; + default: + e = hipErrorInvalidValue; + break; + } + } + + unsigned supportedFlags = hipDeviceScheduleMask | hipDeviceMapHost | hipDeviceLmemResizeToMax; + + if (flags & ~supportedFlags) { + e = hipErrorInvalidValue; + } } else { e = hipErrorInvalidDevice; } diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp index bdb0f869ba..2d46c1ead8 100644 --- a/projects/clr/hipamd/src/hip_hcc.cpp +++ b/projects/clr/hipamd/src/hip_hcc.cpp @@ -66,7 +66,7 @@ int HIP_ATP_MARKER= 0; int HIP_DB= 0; int HIP_VISIBLE_DEVICES = 0; /* Contains a comma-separated sequence of GPU identifiers */ int HIP_NUM_KERNELS_INFLIGHT = 128; -int HIP_BLOCKING_SYNC = 0; +int HIP_FORCE_BLOCKING_SYNC = 0; #define HIP_USE_PRODUCT_NAME 0 //#define DISABLE_COPY_EXT 1 @@ -82,6 +82,7 @@ bool g_visible_device = false; unsigned g_deviceCnt; std::vector g_hip_visible_devices; hsa_agent_t g_cpu_agent; +unsigned g_numLogicalThreads; /* Implementation of malloc and free device functions. @@ -240,6 +241,19 @@ ihipStream_t::ihipStream_t(ihipCtx_t *ctx, hc::accelerator_view av, unsigned int _ctx(ctx), _criticalData(av) { + unsigned schedBits = ctx->_ctxFlags & hipDeviceScheduleMask; + + switch (schedBits) { + case hipDeviceScheduleAuto : _scheduleMode = Auto; break; + case hipDeviceScheduleSpin : _scheduleMode = Spin; break; + case hipDeviceScheduleYield : _scheduleMode = Yield; break; + case hipDeviceScheduleBlockingSync : _scheduleMode = Yield; break; + default:_scheduleMode = Auto; + }; + + + + tprintf(DB_SYNC, " streamCreate: stream=%p\n", this); }; @@ -256,7 +270,23 @@ void ihipStream_t::wait(LockedAccessor_StreamCrit_t &crit, bool assertQueueEmpty { if (! assertQueueEmpty) { tprintf (DB_SYNC, "stream %p wait for queue-empty..\n", this); - crit->_av.wait(HIP_BLOCKING_SYNC ? hc::hcWaitModeBlocked : hc::hcWaitModeActive); + hc::hcWaitMode waitMode = hc::hcWaitModeActive; + + if (_scheduleMode == Auto) { + if (g_deviceCnt > g_numLogicalThreads) { + waitMode = hc::hcWaitModeActive; + } else { + waitMode = hc::hcWaitModeBlocked; + } + } else if (_scheduleMode == Spin) { + waitMode = hc::hcWaitModeActive; + } else if (_scheduleMode == Yield) { + waitMode = hc::hcWaitModeBlocked; + } else { + assert(0); // bad wait mode. + } + + crit->_av.wait(HIP_FORCE_BLOCKING_SYNC ? hc::hcWaitModeBlocked : waitMode); } crit->_kernelCnt = 0; @@ -1090,7 +1120,7 @@ void ihipInit() READ_ENV_I(release, HIP_VISIBLE_DEVICES, CUDA_VISIBLE_DEVICES, "Only devices whose index is present in the secquence are visible to HIP applications and they are enumerated in the order of secquence" ); - READ_ENV_I(release, HIP_BLOCKING_SYNC, 0, "Use blocking synchronization for stream waits. This may increase latency but is friendlier to other processes. If 0, spin-wait."); + READ_ENV_I(release, HIP_FORCE_BLOCKING_SYNC, 0, "Force blocking synchronization for stream waits. This may increase latency but is friendlier to other processes. If 0, used ."); READ_ENV_I(release, HIP_NUM_KERNELS_INFLIGHT, 128, "Max number of inflight kernels per stream before active synchronization is forced."); @@ -1174,13 +1204,14 @@ void ihipInit() g_deviceCnt++; } } + g_numLogicalThreads = std::thread::hardware_concurrency(); // If HIP_VISIBLE_DEVICES is not set, make sure all devices are initialized if(!g_visible_device) { assert(deviceCnt == g_deviceCnt); } - tprintf(DB_SYNC, "pid=%u %-30s\n", getpid(), ""); + tprintf(DB_SYNC, "pid=%u %-30s g_numLogicalThreads=%u\n", getpid(), "", g_numLogicalThreads); } diff --git a/projects/clr/hipamd/src/hip_hcc.h b/projects/clr/hipamd/src/hip_hcc.h index d8045d4cb7..4e05450b19 100644 --- a/projects/clr/hipamd/src/hip_hcc.h +++ b/projects/clr/hipamd/src/hip_hcc.h @@ -384,11 +384,14 @@ typedef ihipStreamCriticalBase_t ihipStreamCritical_t; typedef LockedAccessor LockedAccessor_StreamCrit_t; + //--- // Internal stream structure. class ihipStream_t { public: -typedef uint64_t SeqNum_t ; + enum ScheduleMode {Auto, Spin, Yield}; + typedef uint64_t SeqNum_t ; + ihipStream_t(ihipCtx_t *ctx, hc::accelerator_view av, unsigned int flags); ~ihipStream_t(); @@ -457,6 +460,8 @@ private: // Data // Friends: friend std::ostream& operator<<(std::ostream& os, const ihipStream_t& s); friend hipError_t hipStreamQuery(hipStream_t); + + ScheduleMode _scheduleMode; };