AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

P4 to Git Change 1809277 by gandryey@gera-win10 on 2019/06/11 17:34:13

Browse Source 
SWDEV-180872 - Runtime support changes for Cooperative Group Features
	- Initial implementation of the core functionality. Disabled by default. Use GPU_ENABLE_COOP_GROUPS=1 to enable the feature.
	- Runtime uses device queue for cooperative executions with a synchronization on the launched queue.
	- The current implementation is pure runtime change and it can work if only one app uses this feature. No ROCr/KFD support was added or tested
	- Only inline assembler was tested

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_device.cpp#20 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_device_runtime.cpp#15 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.def.in#15 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.map.in#17 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#28 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#32 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/device.hpp#338 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpudevice.cpp#606 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpudevice.hpp#171 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palblit.cpp#31 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palblit.hpp#9 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.cpp#142 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.hpp#39 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palschedcl.cpp#6 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.cpp#135 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.hpp#61 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocblit.cpp#32 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocblit.hpp#12 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.cpp#127 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.hpp#37 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocschedcl.cpp#3 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#75 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.hpp#23 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/command.cpp#94 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/command.hpp#92 edit
... //depot/stg/opencl/drivers/opencl/runtime/utils/flags.hpp#311 edit
This commit is contained in:
foreman
2019-06-12 10:00:38 -04:00
parent d1abc0baa2
commit 517bf08c37
20 changed files with 437 additions and 47 deletions
Download Patch File Download Diff File Expand all files Collapse all files
rocclr/runtime/device/device.hpp
+8 -1
View File
@@ -507,6 +507,11 @@ struct Info : public amd::EmbeddedObject {
//! Max numbers of threads per CU
cl_uint maxThreadsPerCU_;
//! GPU device supports a launch of cooperative groups
cl_bool cooperativeGroups_;
//! GPU device supports a launch of cooperative groups on multiple devices
cl_bool cooperativeMultiDeviceGroups_;
};
//! Device settings
@@ -1286,7 +1291,9 @@ class Device : public RuntimeObject {
virtual void svmFree(void* ptr) const = 0;
//! Validate kernel
virtual bool validateKernel(const amd::Kernel& kernel, const device::VirtualDevice* vdev) {
virtual bool validateKernel(const amd::Kernel& kernel,
const device::VirtualDevice* vdev,
bool coop_groups = false) {
return true;
};
rocclr/runtime/device/gpu/gpudevice.cpp
+2 -1
View File
@@ -1955,7 +1955,8 @@ bool Device::allocScratch(uint regNum, const VirtualGPU* vgpu) {
return true;
}
bool Device::validateKernel(const amd::Kernel& kernel, const device::VirtualDevice* vdev) {
bool Device::validateKernel(
const amd::Kernel& kernel, const device::VirtualDevice* vdev, bool coop_groups) {
// Find the number of scratch registers used in the kernel
const device::Kernel* devKernel = kernel.getDeviceKernel(*this);
uint regNum = static_cast<uint>(devKernel->workGroupInfo()->scratchRegs_);
rocclr/runtime/device/gpu/gpudevice.hpp
+2 -1
View File
@@ -410,7 +410,8 @@ class Device : public NullDevice, public CALGSLDevice {
//! Validates kernel before execution
virtual bool validateKernel(const amd::Kernel& kernel, //!< AMD kernel object
const device::VirtualDevice* vdev);
const device::VirtualDevice* vdev,
bool coop_groups = false);
virtual bool SetClockMode(const cl_set_device_clock_mode_input_amd setClockModeInput, cl_set_device_clock_mode_output_amd* pSetClockModeOutput);
rocclr/runtime/device/pal/palblit.cpp
+23
View File
@@ -2381,6 +2381,29 @@ void KernelBlitManager::writeRawData(device::Memory& memory, size_t size, const
synchronize();
}
bool KernelBlitManager::RunGwsInit(
uint32_t value) const {
amd::ScopedLock k(lockXferOps_);
size_t globalWorkOffset[1] = { 0 };
size_t globalWorkSize[1] = { 1 };
size_t localWorkSize[1] = { 1 };
// Program kernels arguments
setArgument(kernels_[GwsInit], 0, sizeof(uint32_t), &value);
// Create ND range object for the kernel's execution
amd::NDRangeContainer ndrange(1, globalWorkOffset, globalWorkSize, localWorkSize);
// Execute the blit
address parameters = kernels_[GwsInit]->parameters().values();
bool result = gpu().submitKernelInternal(ndrange, *kernels_[GwsInit], parameters);
synchronize();
return result;
}
amd::Memory* DmaBlitManager::pinHostMemory(const void* hostMem, size_t pinSize,
size_t& partial) const {
size_t pinAllocSize;
rocclr/runtime/device/pal/palblit.hpp
+6 -1
View File
@@ -205,6 +205,7 @@ class KernelBlitManager : public DmaBlitManager {
FillBuffer,
FillImage,
Scheduler,
GwsInit,
BlitTotal
};
@@ -352,6 +353,10 @@ class KernelBlitManager : public DmaBlitManager {
const void* data //!< Raw data pointer
) const;
//! Runs a blit kernel for GWS init
bool RunGwsInit(uint32_t value //!< Initial value for GWS resource
) const;
virtual amd::Monitor* lockXfer() const { return &lockXferOps_; }
private:
@@ -406,7 +411,7 @@ static const char* BlitName[KernelBlitManager::BlitTotal] = {
"copyImage", "copyImage1DA", "copyImageToBuffer",
"copyBufferToImage", "copyBufferRect", "copyBufferRectAligned",
"copyBuffer", "copyBufferAligned", "fillBuffer",
"fillImage", "scheduler",
"fillImage", "scheduler", "gwsInit"
};
/*@}*/ // namespace pal
rocclr/runtime/device/pal/paldevice.cpp
+48 -4
View File
@@ -660,6 +660,8 @@ void NullDevice::fillDeviceInfo(const Pal::DeviceProperties& palProp,
info_.pcieRevisionId_ = palProp.revisionId;
info_.maxThreadsPerCU_ = info_.wavefrontWidth_ * info_.simdPerCU_ *
palProp.gfxipProperties.shaderCore.numWavefrontsPerSimd;
info_.cooperativeGroups_ = GPU_ENABLE_COOP_GROUPS;
info_.cooperativeMultiDeviceGroups_ = GPU_ENABLE_COOP_GROUPS;
}
}
@@ -852,6 +854,7 @@ Device::~Device() {
}
extern const char* SchedulerSourceCode;
extern const char* GwsInitSourceCode;
Pal::IDevice* gDeviceList[Pal::MaxDevices] = {};
uint32_t gStartDevice = 0;
uint32_t gNumDevices = 0;
@@ -2059,7 +2062,8 @@ bool Device::allocScratch(uint regNum, const VirtualGPU* vgpu, uint vgprs) {
return true;
}
bool Device::validateKernel(const amd::Kernel& kernel, const device::VirtualDevice* vdev) {
bool Device::validateKernel(
const amd::Kernel& kernel, const device::VirtualDevice* vdev, bool coop_groups) {
// Find the number of scratch registers used in the kernel
const device::Kernel* devKernel = kernel.getDeviceKernel(*this);
uint regNum = static_cast<uint>(devKernel->workGroupInfo()->scratchRegs_);
@@ -2068,6 +2072,14 @@ bool Device::validateKernel(const amd::Kernel& kernel, const device::VirtualDevi
if (!allocScratch(regNum, vgpu, devKernel->workGroupInfo()->usedVGPRs_)) {
return false;
}
// Runtime plans to launch cooperative groups on the device queue, thus
// validate the scratch buffer on that queue
if (coop_groups) {
vgpu = xferQueue();
if (!allocScratch(regNum, vgpu, devKernel->workGroupInfo()->usedVGPRs_)) {
return false;
}
}
if (devKernel->hsa()) {
const HSAILKernel* hsaKernel = static_cast<const HSAILKernel*>(devKernel);
@@ -2224,6 +2236,35 @@ void Device::svmFree(void* ptr) const {
}
}
bool Device::AcquireExclusiveGpuAccess() {
// Lock the virtual GPU list
vgpusAccess().lock();
// Find all available virtual GPUs and lock them
// from the execution of commands
for (uint idx = 0; idx < vgpus().size(); ++idx) {
vgpus()[idx]->execution().lock();
// Make sure a wait is done
vgpus()[idx]->WaitForIdleCompute();
}
// if (!hsa_exclusive_gpu_access_) {
// @todo call rocr
// hsa_exclusive_gpu_access_ = true;
// }
return true;
}
void Device::ReleaseExclusiveGpuAccess(VirtualGPU& vgpu) const {
vgpu.WaitForIdleCompute();
// Find all available virtual GPUs and unlock them
// for the execution of commands
for (uint idx = 0; idx < vgpus().size(); ++idx) {
vgpus()[idx]->execution().unlock();
}
// Unock the virtual GPU list
vgpusAccess().unlock();
}
Device::SrdManager::~SrdManager() {
for (uint i = 0; i < pool_.size(); ++i) {
@@ -2333,7 +2374,7 @@ bool Device::createBlitProgram() {
bool result = true;
// Delayed compilation due to brig_loader memory allocation
const char* scheduler = nullptr;
const char* blits = nullptr;
const char* ocl20 = nullptr;
std::string sch = SchedulerSourceCode;
@@ -2348,14 +2389,17 @@ bool Device::createBlitProgram() {
sch.replace(loc, sizeof(AmdScheduler) - 1, AmdSchedulerPal);
loc = sch.find(AmdScheduler, (loc + sizeof(AmdSchedulerPal) - 1));
sch.replace(loc, sizeof(AmdScheduler) - 1, AmdSchedulerPal);
if (info().cooperativeGroups_) {
sch.append(GwsInitSourceCode);
}
}
scheduler = sch.c_str();
blits = sch.c_str();
ocl20 = "-cl-std=CL2.0";
}
blitProgram_ = new BlitProgram(context_);
// Create blit programs
if (blitProgram_ == nullptr || !blitProgram_->create(this, scheduler, ocl20)) {
if (blitProgram_ == nullptr || !blitProgram_->create(this, blits, ocl20)) {
delete blitProgram_;
blitProgram_ = nullptr;
LogError("Couldn't create blit kernels!");
rocclr/runtime/device/pal/paldevice.hpp
+5 -1
View File
@@ -342,7 +342,8 @@ class Device : public NullDevice {
//! Validates kernel before execution
virtual bool validateKernel(const amd::Kernel& kernel, //!< AMD kernel object
const device::VirtualDevice* vdev);
const device::VirtualDevice* vdev,
bool coop_group = false);
virtual bool SetClockMode(const cl_set_device_clock_mode_input_amd setClockModeInput,
cl_set_device_clock_mode_output_amd* pSetClockModeOutput);
@@ -525,6 +526,9 @@ class Device : public NullDevice {
}
}
bool AcquireExclusiveGpuAccess();
void ReleaseExclusiveGpuAccess(VirtualGPU& vgpu) const;
private:
static void PAL_STDCALL PalDeveloperCallback(void* pPrivateData, const Pal::uint32 deviceIndex,
Pal::Developer::CallbackType type, void* pCbData);
rocclr/runtime/device/pal/palschedcl.cpp
+22 -2
View File
@@ -3,9 +3,9 @@
//
namespace pal {
#define SCHEDULER_KERNEL(...) #__VA_ARGS__
#define BLIT_KERNEL(...) #__VA_ARGS__
const char* SchedulerSourceCode = SCHEDULER_KERNEL(
const char* SchedulerSourceCode = BLIT_KERNEL(
%s
\n
extern void __amd_scheduler(__global void*, __global void*, uint);
@@ -15,4 +15,24 @@ __kernel void scheduler(__global void* queue, __global void* params, uint paramI
}
\n);
const char* GwsInitSourceCode = BLIT_KERNEL(
\n
__kernel void gwsInit(uint value) {
unsigned int m0_backup, new_m0;
__asm__ __volatile__(
"s_mov_b32 %0 m0\n"
"v_readfirstlane_b32 %1 %2\n"
"s_nop 0\n"
"s_mov_b32 m0 %1\n"
"s_nop 0\n"
"ds_gws_init %3 offset:0 gds\n"
"s_waitcnt lgkmcnt(0) expcnt(0)\n"
"s_mov_b32 m0 %0\n"
"s_nop 0"
: "=s"(m0_backup), "=s"(new_m0)
: "v"(0 << 0x10), "{v0}"(value - 1)
: "memory");
}
\n);
} // namespace pal
rocclr/runtime/device/pal/palvirtual.cpp
+50 -10
View File
@@ -2166,24 +2166,62 @@ void VirtualGPU::PostDeviceEnqueue(const amd::Kernel& kernel, const HSAILKernel&
// ================================================================================================
void VirtualGPU::submitKernel(amd::NDRangeKernelCommand& vcmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
if (vcmd.cooperativeGroups()) {
uint32_t workgroups = 0;
for (uint i = 0; i < vcmd.sizes().dimensions(); i++) {
if ((vcmd.sizes().local()[i] != 0) && (vcmd.sizes().global()[i] != 1)) {
workgroups += (vcmd.sizes().global()[i] / vcmd.sizes().local()[i]);
}
}
uint32_t counter = workgroups *
amd::alignUp(vcmd.sizes().local().product(), dev().info().wavefrontWidth_) /
dev().info().wavefrontWidth_;
profilingBegin(vcmd);
bool test = true;
VirtualGPU* queue = (test) ? this : dev().xferQueue();
// Submit kernel to HW
if (!submitKernelInternal(vcmd.sizes(), vcmd.kernel(), vcmd.parameters(), false, &vcmd.event(),
vcmd.sharedMemBytes())) {
vcmd.setStatus(CL_INVALID_OPERATION);
// Wait for the execution on the current queue, since the coop groups will use the device queue
waitAllEngines();
amd::ScopedLock lock(queue->blitMgr().lockXfer());
queue->profilingBegin(vcmd);
static_cast<KernelBlitManager&>(queue->blitMgr()).RunGwsInit(counter);
queue->addBarrier(RgpSqqtBarrierReason::PostDeviceEnqueue);
// Submit kernel to HW
if (!queue->submitKernelInternal(vcmd.sizes(), vcmd.kernel(), vcmd.parameters(),
false, &vcmd.event(), vcmd.sharedMemBytes(),
vcmd.cooperativeGroups())) {
vcmd.setStatus(CL_INVALID_OPERATION);
}
queue->profilingEnd(vcmd);
// Wait for the execution on the device queue. Keep the current queue in-order
queue->waitAllEngines();
} else {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
profilingBegin(vcmd);
// Submit kernel to HW
if (!submitKernelInternal(vcmd.sizes(), vcmd.kernel(), vcmd.parameters(), false, &vcmd.event(),
vcmd.sharedMemBytes(), vcmd.cooperativeGroups())) {
vcmd.setStatus(CL_INVALID_OPERATION);
}
profilingEnd(vcmd);
}
profilingEnd(vcmd);
}
// ================================================================================================
bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const amd::Kernel& kernel,
const_address parameters, bool nativeMem,
amd::Event* enqueueEvent, uint32_t sharedMemBytes) {
amd::Event* enqueueEvent, uint32_t sharedMemBytes,
bool cooperativeGroup) {
size_t newOffset[3] = {0, 0, 0};
size_t newGlobalSize[3] = {0, 0, 0};
@@ -2239,6 +2277,7 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const
LogError("Wrong memory objects!");
return false;
}
bool needFlush = false;
// Avoid flushing when PerfCounter is enabled, to make sure PerfStart/dispatch/PerfEnd
// are in the same cmdBuffer
@@ -2358,6 +2397,7 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const
return false;
}
}
// Perform post dispatch logic for RGP traces
if (rgpCaptureEna()) {
dev().rgpCaptureMgr()->PostDispatch(this);
rocclr/runtime/device/pal/palvirtual.hpp
+10 -1
View File
@@ -298,7 +298,8 @@ class VirtualGPU : public device::VirtualDevice {
const_address parameters, //!< Parameters for the kernel
bool nativeMem = true, //!< Native memory objects
amd::Event* enqueueEvent = nullptr, //!< Event provided in the enqueue kernel command
uint32_t sharedMemBytes = 0 //!< Shared memory size
uint32_t sharedMemBytes = 0, //!< Shared memory size
bool cooperativeGroups = false //!< TRUE if cooperative groups mode is required
);
void submitNativeFn(amd::NativeFnCommand& vcmd);
void submitFillMemory(amd::FillMemoryCommand& vcmd);
@@ -512,6 +513,14 @@ class VirtualGPU : public device::VirtualDevice {
//! Checks if RGP capture is enabled
bool rgpCaptureEna() const { return state_.rgpCaptureEnabled_; }
//! Waits for idle on compute engine
void WaitForIdleCompute() {
if (events_[MainEngine].isValid()) {
queues_[events_[MainEngine].engineId_]->waitForEvent(events_[MainEngine].id_);
events_[MainEngine].invalidate();
}
}
protected:
void profileEvent(EngineType engine, bool type) const;
rocclr/runtime/device/rocm/rocblit.cpp
+24
View File
@@ -2235,4 +2235,28 @@ bool KernelBlitManager::runScheduler(uint64_t vqVM, amd::Memory* schedulerParam,
return true;
}
bool KernelBlitManager::RunGwsInit(
uint32_t value) const {
amd::ScopedLock k(lockXferOps_);
size_t globalWorkOffset[1] = { 0 };
size_t globalWorkSize[1] = { 1 };
size_t localWorkSize[1] = { 1 };
// Program kernels arguments
setArgument(kernels_[GwsInit], 0, sizeof(uint32_t), &value);
// Create ND range object for the kernel's execution
amd::NDRangeContainer ndrange(1, globalWorkOffset, globalWorkSize, localWorkSize);
// Execute the blit
address parameters = captureArguments(kernels_[GwsInit]);
bool result = gpu().submitKernelInternal(ndrange, *kernels_[GwsInit], parameters, nullptr);
releaseArguments(parameters);
return result;
}
} // namespace pal
rocclr/runtime/device/rocm/rocblit.hpp
+6 -1
View File
@@ -234,6 +234,7 @@ class KernelBlitManager : public DmaBlitManager {
FillBuffer,
FillImage,
Scheduler,
GwsInit,
BlitTotal
};
@@ -374,6 +375,10 @@ class KernelBlitManager : public DmaBlitManager {
hsa_signal_t& schedulerSignal,
uint threads);
//! Runs a blit kernel for GWS init
bool RunGwsInit(uint32_t value //!< Initial value for GWS resource
) const;
private:
static const size_t MaxXferBuffers = 2;
static const uint TransferSplitSize = 1;
@@ -434,7 +439,7 @@ static const char* BlitName[KernelBlitManager::BlitTotal] = {
"copyImage", "copyImage1DA", "copyImageToBuffer",
"copyBufferToImage", "copyBufferRect", "copyBufferRectAligned",
"copyBuffer", "copyBufferAligned", "fillBuffer",
"fillImage", "scheduler",
"fillImage", "scheduler", "gwsInit"
};
inline void KernelBlitManager::setArgument(amd::Kernel* kernel, size_t index, size_t size, const void* value) const {
rocclr/runtime/device/rocm/rocdevice.cpp
+43
View File
@@ -148,6 +148,7 @@ Device::Device(hsa_agent_t bkendDevice)
, pro_device_(nullptr)
, pro_ena_(false)
, freeMem_(0)
, hsa_exclusive_gpu_access_(false)
, numOfVgpus_(0) {
group_segment_.handle = 0;
system_segment_.handle = 0;
@@ -585,6 +586,7 @@ bool Device::init() {
}
extern const char* SchedulerSourceCode;
extern const char* GwsInitSourceCode;
void Device::tearDown() {
NullDevice::tearDown();
@@ -648,6 +650,9 @@ bool Device::create(bool sramEccEnabled) {
#if defined(WITH_LIGHTNING_COMPILER) || defined(USE_COMGR_LIBRARY)
std::string sch = SchedulerSourceCode;
if (settings().useLightning_) {
if (info().cooperativeGroups_) {
sch.append(GwsInitSourceCode);
}
scheduler = sch.c_str();
}
#ifndef USE_COMGR_LIBRARY
@@ -781,6 +786,38 @@ device::Program* NullDevice::createProgram(amd::option::Options* options) {
return program;
}
bool Device::AcquireExclusiveGpuAccess() {
// Lock the virtual GPU list
vgpusAccess().lock();
// Find all available virtual GPUs and lock them
// from the execution of commands
for (uint idx = 0; idx < vgpus().size(); ++idx) {
vgpus()[idx]->execution().lock();
// Make sure a wait is done
vgpus()[idx]->releaseGpuMemoryFence();
}
if (!hsa_exclusive_gpu_access_) {
// @todo call rocr
hsa_exclusive_gpu_access_ = true;
}
return true;
}
void Device::ReleaseExclusiveGpuAccess(VirtualGPU& vgpu) const {
// Make sure the operation is done
vgpu.releaseGpuMemoryFence();
// Find all available virtual GPUs and unlock them
// for the execution of commands
for (uint idx = 0; idx < vgpus().size(); ++idx) {
vgpus()[idx]->execution().unlock();
}
// Unock the virtual GPU list
vgpusAccess().unlock();
}
device::Program* Device::createProgram(amd::option::Options* options) {
device::Program* program;
if (settings().useLightning_) {
@@ -1339,6 +1376,8 @@ bool Device::populateOCLDeviceConstants() {
//TODO: set to true once thread trace support is available
info_.threadTraceEnable_ = false;
info_.pcieDeviceId_ = deviceInfo_.pciDeviceId_;
info_.cooperativeGroups_ = GPU_ENABLE_COOP_GROUPS;
info_.cooperativeMultiDeviceGroups_ = GPU_ENABLE_COOP_GROUPS;
}
info_.maxPipePacketSize_ = info_.maxMemAllocSize_;
@@ -1356,8 +1395,12 @@ bool Device::populateOCLDeviceConstants() {
}
device::VirtualDevice* Device::createVirtualDevice(amd::CommandQueue* queue) {
amd::ScopedLock lock(vgpusAccess());
bool profiling = (queue != nullptr) && queue->properties().test(CL_QUEUE_PROFILING_ENABLE);
profiling |= (queue == nullptr) ? true : false;
// Initialization of heap and other resources occur during the command
// queue creation time.
VirtualGPU* virtualDevice = new VirtualGPU(*this);
rocclr/runtime/device/rocm/rocdevice.hpp
+15 -2
View File
@@ -393,6 +393,19 @@ class Device : public NullDevice {
virtual amd::Memory* IpcAttach(const void* handle, size_t mem_size, unsigned int flags, void** dev_ptr) const;
virtual void IpcDetach (amd::Memory& memory) const;
bool AcquireExclusiveGpuAccess();
void ReleaseExclusiveGpuAccess(VirtualGPU& vgpu) const;
//! Returns the lock object for the virtual gpus list
amd::Monitor& vgpusAccess() const { return vgpusAccess_; }
typedef std::vector<VirtualGPU*> VirtualGPUs;
//! Returns the list of all virtual GPUs running on this device
const VirtualGPUs& vgpus() const { return vgpus_; }
VirtualGPUs vgpus_; //!< The list of all running virtual gpus (lock protected)
VirtualGPU* xferQueue() const;
private:
static hsa_ven_amd_loader_1_00_pfn_t amd_loader_ext_table;
@@ -416,13 +429,13 @@ class Device : public NullDevice {
amd::Context* context_; //!< A dummy context for internal data transfer
VirtualGPU* xferQueue_; //!< Transfer queue, created on demand
VirtualGPU* xferQueue() const;
XferBuffers* xferRead_; //!< Transfer buffers read
XferBuffers* xferWrite_; //!< Transfer buffers write
const IProDevice* pro_device_; //!< AMDGPUPro device
bool pro_ena_; //!< Extra functionality with AMDGPUPro device, beyond ROCr
std::atomic<size_t> freeMem_; //!< Total of free memory available
mutable amd::Monitor vgpusAccess_; //!< Lock to serialise virtual gpu list access
bool hsa_exclusive_gpu_access_; //!< TRUE if current device was moved into exclusive GPU access mode
public:
amd::Atomic<uint> numOfVgpus_; //!< Virtual gpu unique index
rocclr/runtime/device/rocm/rocschedcl.cpp
+21 -2
View File
@@ -3,9 +3,9 @@
//
namespace roc {
#define SCHEDULER_KERNEL(...) #__VA_ARGS__
#define BLIT_KERNEL(...) #__VA_ARGS__
const char* SchedulerSourceCode = SCHEDULER_KERNEL(
const char* SchedulerSourceCode = BLIT_KERNEL(
\n
extern void __amd_scheduler_rocm(__global void*);
\n
@@ -13,5 +13,24 @@ __kernel void scheduler(__global void* params) {
__amd_scheduler_rocm(params);
}
\n);
const char* GwsInitSourceCode = BLIT_KERNEL(
\n
__kernel void gwsInit(uint value) {
unsigned int m0_backup, new_m0;
__asm__ __volatile__(
"s_mov_b32 %0 m0\n"
"v_readfirstlane_b32 %1 %2\n"
"s_nop 0\n"
"s_mov_b32 m0 %1\n"
"s_nop 0\n"
"ds_gws_init %3 offset:0 gds\n"
"s_waitcnt lgkmcnt(0) expcnt(0)\n"
"s_mov_b32 m0 %0\n"
"s_nop 0"
: "=s"(m0_backup), "=s"(new_m0)
: "v"(0 << 0x10), "{v0}"(value - 1)
: "memory");
}
\n);
} // namespace roc
rocclr/runtime/device/rocm/rocvirtual.cpp
+118 -13
View File
@@ -186,13 +186,16 @@ void VirtualGPU::MemoryDependency::clear(bool all) {
}
}
bool VirtualGPU::processMemObjects(const amd::Kernel& kernel, const_address params, size_t& ldsAddress) {
bool VirtualGPU::processMemObjects(const amd::Kernel& kernel, const_address params,
size_t& ldsAddress, bool cooperativeGroups) {
Kernel& hsaKernel = const_cast<Kernel&>(static_cast<const Kernel&>(*(kernel.getDeviceKernel(dev()))));
const amd::KernelSignature& signature = kernel.signature();
const amd::KernelParameters& kernelParams = kernel.parameters();
// AQL packets
setAqlHeader(kDispatchPacketHeaderNoSync);
if (!cooperativeGroups) {
// AQL packets
setAqlHeader(kDispatchPacketHeaderNoSync);
}
// Mark the tracker with a new kernel,
// so we can avoid checks of the aliased objects
@@ -531,6 +534,7 @@ bool VirtualGPU::releaseGpuMemoryFence() {
VirtualGPU::VirtualGPU(Device& device)
: device::VirtualDevice(device),
execution_("Virtual GPU execution lock", true),
roc_device_(device),
virtualQueue_(nullptr),
deviceQueueSize_(0),
@@ -553,6 +557,10 @@ VirtualGPU::VirtualGPU(Device& device)
kernarg_pool_cur_offset_ = 0;
aqlHeader_ = kDispatchPacketHeaderNoSync;
barrier_signal_.handle = 0;
// Note: Virtual GPU device creation must be a thread safe operation
roc_device_.vgpus_.resize(roc_device_.numOfVgpus_);
roc_device_.vgpus_[index()] = this;
}
VirtualGPU::~VirtualGPU() {
@@ -598,6 +606,10 @@ VirtualGPU::~VirtualGPU() {
}
--roc_device_.numOfVgpus_; // Virtual gpu unique index decrementing
roc_device_.vgpus_.erase(roc_device_.vgpus_.begin() + index());
for (uint idx = index(); idx < roc_device_.vgpus().size(); ++idx) {
roc_device_.vgpus()[idx]->index_--;
}
}
bool VirtualGPU::create(bool profilingEna) {
@@ -844,6 +856,9 @@ void VirtualGPU::updateCommandsState(amd::Command* list) {
}
void VirtualGPU::submitReadMemory(amd::ReadMemoryCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
// Wait on a kernel if one is outstanding
releaseGpuMemoryFence();
@@ -932,6 +947,9 @@ void VirtualGPU::submitReadMemory(amd::ReadMemoryCommand& cmd) {
}
void VirtualGPU::submitWriteMemory(amd::WriteMemoryCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
// Wait on a kernel if one is outstanding
releaseGpuMemoryFence();
@@ -1026,6 +1044,9 @@ void VirtualGPU::submitWriteMemory(amd::WriteMemoryCommand& cmd) {
}
void VirtualGPU::submitSvmFreeMemory(amd::SvmFreeMemoryCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
// in-order semantics: previous commands need to be done before we start
releaseGpuMemoryFence();
@@ -1125,6 +1146,9 @@ bool VirtualGPU::copyMemory(cl_command_type type, amd::Memory& srcMem, amd::Memo
}
void VirtualGPU::submitCopyMemory(amd::CopyMemoryCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
// Wait on a kernel if one is outstanding
releaseGpuMemoryFence();
@@ -1142,6 +1166,9 @@ void VirtualGPU::submitCopyMemory(amd::CopyMemoryCommand& cmd) {
}
void VirtualGPU::submitSvmCopyMemory(amd::SvmCopyMemoryCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
// in-order semantics: previous commands need to be done before we start
releaseGpuMemoryFence();
@@ -1218,6 +1245,9 @@ void VirtualGPU::submitSvmCopyMemory(amd::SvmCopyMemoryCommand& cmd) {
}
void VirtualGPU::submitCopyMemoryP2P(amd::CopyMemoryP2PCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
// Wait on a kernel if one is outstanding
releaseGpuMemoryFence();
@@ -1303,6 +1333,9 @@ void VirtualGPU::submitCopyMemoryP2P(amd::CopyMemoryP2PCommand& cmd) {
}
void VirtualGPU::submitSvmMapMemory(amd::SvmMapMemoryCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
// Wait on a kernel if one is outstanding
releaseGpuMemoryFence();
@@ -1339,6 +1372,9 @@ void VirtualGPU::submitSvmMapMemory(amd::SvmMapMemoryCommand& cmd) {
}
void VirtualGPU::submitSvmUnmapMemory(amd::SvmUnmapMemoryCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
// Wait on a kernel if one is outstanding
releaseGpuMemoryFence();
@@ -1376,6 +1412,9 @@ void VirtualGPU::submitSvmUnmapMemory(amd::SvmUnmapMemoryCommand& cmd) {
}
void VirtualGPU::submitMapMemory(amd::MapMemoryCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
// Wait on a kernel if one is outstanding
releaseGpuMemoryFence();
@@ -1471,6 +1510,9 @@ void VirtualGPU::submitMapMemory(amd::MapMemoryCommand& cmd) {
}
void VirtualGPU::submitUnmapMemory(amd::UnmapMemoryCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
roc::Memory* devMemory = static_cast<roc::Memory*>(cmd.memory().getDeviceMemory(dev(), false));
const device::Memory::WriteMapInfo* mapInfo = devMemory->writeMapInfo(cmd.mapPtr());
@@ -1560,6 +1602,9 @@ void VirtualGPU::submitUnmapMemory(amd::UnmapMemoryCommand& cmd) {
bool VirtualGPU::fillMemory(cl_command_type type, amd::Memory* amdMemory, const void* pattern,
size_t patternSize, const amd::Coord3D& origin,
const amd::Coord3D& size) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
Memory* memory = dev().getRocMemory(amdMemory);
bool entire = amdMemory->isEntirelyCovered(origin, size);
@@ -1615,6 +1660,9 @@ bool VirtualGPU::fillMemory(cl_command_type type, amd::Memory* amdMemory, const
}
void VirtualGPU::submitFillMemory(amd::FillMemoryCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
// Wait on a kernel if one is outstanding
releaseGpuMemoryFence();
@@ -1628,6 +1676,9 @@ void VirtualGPU::submitFillMemory(amd::FillMemoryCommand& cmd) {
}
void VirtualGPU::submitSvmFillMemory(amd::SvmFillMemoryCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
// in-order semantics: previous commands need to be done before we start
releaseGpuMemoryFence();
@@ -1669,6 +1720,9 @@ void VirtualGPU::submitSvmFillMemory(amd::SvmFillMemoryCommand& cmd) {
}
void VirtualGPU::submitMigrateMemObjects(amd::MigrateMemObjectsCommand& vcmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
// Wait on a kernel if one is outstanding
releaseGpuMemoryFence();
@@ -1875,13 +1929,13 @@ bool VirtualGPU::createVirtualQueue(uint deviceQueueSize)
}
bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const amd::Kernel& kernel,
const_address parameters, void* eventHandle, uint32_t sharedMemBytes) {
const_address parameters, void* eventHandle, uint32_t sharedMemBytes, bool cooperativeGroups) {
device::Kernel* devKernel = const_cast<device::Kernel*>(kernel.getDeviceKernel(dev()));
Kernel& gpuKernel = static_cast<Kernel&>(*devKernel);
size_t ldsUsage = gpuKernel.WorkgroupGroupSegmentByteSize();
// Check memory dependency and SVM objects
if (!processMemObjects(kernel, parameters, ldsUsage)) {
if (!processMemObjects(kernel, parameters, ldsUsage, cooperativeGroups)) {
LogError("Wrong memory objects!");
return false;
}
@@ -2087,16 +2141,57 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const
* the list of kernel parameters.
*/
void VirtualGPU::submitKernel(amd::NDRangeKernelCommand& vcmd) {
profilingBegin(vcmd);
if (vcmd.cooperativeGroups()) {
uint32_t workgroups = 0;
for (uint i = 0; i < vcmd.sizes().dimensions(); i++) {
if ((vcmd.sizes().local()[i] != 0) && (vcmd.sizes().global()[i] != 1)) {
workgroups += (vcmd.sizes().global()[i] / vcmd.sizes().local()[i]);
}
}
uint32_t counter = workgroups *
amd::alignUp(vcmd.sizes().local().product(), dev().info().wavefrontWidth_) /
dev().info().wavefrontWidth_;
// Submit kernel to HW
if (!submitKernelInternal(vcmd.sizes(), vcmd.kernel(), vcmd.parameters(),
static_cast<void*>(as_cl(&vcmd.event())), vcmd.sharedMemBytes())) {
LogError("AQL dispatch failed!");
vcmd.setStatus(CL_INVALID_OPERATION);
// Get device queue for exclusive GPU access
VirtualGPU* queue = dev().xferQueue();
// Wait for the execution on the current queue, since the coop groups will use the device queue
releaseGpuMemoryFence();
// Lock the queue, using the blit manager lock
amd::ScopedLock lock(queue->blitMgr().lockXfer());
queue->profilingBegin(vcmd);
static_cast<KernelBlitManager&>(queue->blitMgr()).RunGwsInit(counter);
// Sync AQL packets
queue->setAqlHeader(kDispatchPacketHeader);
// Submit kernel to HW
if (!queue->submitKernelInternal(vcmd.sizes(), vcmd.kernel(), vcmd.parameters(),
static_cast<void*>(as_cl(&vcmd.event())), vcmd.sharedMemBytes(), vcmd.cooperativeGroups())) {
LogError("AQL dispatch failed!");
vcmd.setStatus(CL_INVALID_OPERATION);
}
// Wait for the execution on the device queue. Keep the current queue in-order
queue->releaseGpuMemoryFence();
queue->profilingEnd(vcmd);
} else {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
profilingBegin(vcmd);
// Submit kernel to HW
if (!submitKernelInternal(vcmd.sizes(), vcmd.kernel(), vcmd.parameters(),
static_cast<void*>(as_cl(&vcmd.event())), vcmd.sharedMemBytes(), vcmd.cooperativeGroups())) {
LogError("AQL dispatch failed!");
vcmd.setStatus(CL_INVALID_OPERATION);
}
profilingEnd(vcmd);
}
profilingEnd(vcmd);
}
void VirtualGPU::submitNativeFn(amd::NativeFnCommand& cmd) {
@@ -2108,6 +2203,9 @@ void VirtualGPU::submitMarker(amd::Marker& cmd) {
}
void VirtualGPU::submitAcquireExtObjects(amd::AcquireExtObjectsCommand& vcmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
profilingBegin(vcmd);
auto fence = kBarrierAcquirePacket;
dispatchAqlPacket(&fence, false);
@@ -2115,6 +2213,8 @@ void VirtualGPU::submitAcquireExtObjects(amd::AcquireExtObjectsCommand& vcmd) {
}
void VirtualGPU::submitReleaseExtObjects(amd::ReleaseExtObjectsCommand& vcmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
profilingBegin(vcmd);
auto fence = kBarrierReleasePacket;
dispatchAqlPacket(&fence, false);
@@ -2176,6 +2276,9 @@ amd::Memory* VirtualGPU::findPinnedMem(void* addr, size_t size) {
void VirtualGPU::enableSyncBlit() const { blitMgr_->enableSynchronization(); }
void VirtualGPU::submitTransferBufferFromFile(amd::TransferBufferFileCommand& cmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
size_t copySize = cmd.size()[0];
size_t fileOffset = cmd.fileOffset();
Memory* mem = dev().getRocMemory(&cmd.memory());
@@ -2228,6 +2331,8 @@ void VirtualGPU::submitTransferBufferFromFile(amd::TransferBufferFileCommand& cm
}
void VirtualGPU::submitPerfCounter(amd::PerfCounterCommand& vcmd) {
// Make sure VirtualGPU has an exclusive access to the resources
amd::ScopedLock lock(execution());
const amd::PerfCounterCommand::PerfCounterList counters = vcmd.getCounters();
rocclr/runtime/device/rocm/rocvirtual.hpp
+9 -3
View File
@@ -167,7 +167,8 @@ class VirtualGPU : public device::VirtualDevice {
const amd::Kernel& kernel, //!< Kernel for execution
const_address parameters, //!< Parameters for the kernel
void* event_handle, //!< Handle to OCL event for debugging
uint32_t sharedMemBytes = 0 //!< Shared memory size
uint32_t sharedMemBytes = 0, //!< Shared memory size
bool cooperativeGroups = false //!< TRUE if cooperative groups mode is required
);
void submitNativeFn(amd::NativeFnCommand& cmd);
void submitMarker(amd::Marker& cmd);
@@ -220,7 +221,8 @@ class VirtualGPU : public device::VirtualDevice {
//! Detects memory dependency for HSAIL kernels and uses appropriate AQL header
bool processMemObjects(const amd::Kernel& kernel, //!< AMD kernel object for execution
const_address params, //!< Pointer to the param's store
size_t& ldsAddress //!< LDS usage
size_t& ldsAddress, //!< LDS usage
bool cooperativeGroups //!< Dispatch with cooperative groups
);
// Retun the virtual gpu unique index
uint index() const { return index_; }
@@ -242,6 +244,9 @@ class VirtualGPU : public device::VirtualDevice {
void enableSyncBlit() const;
//! Returns the monitor object for execution access by VirtualGPU
amd::Monitor& execution() { return execution_; }
// } roc OpenCL integration
private:
bool dispatchAqlPacket(hsa_kernel_dispatch_packet_t* packet, bool blocking = true);
@@ -296,6 +301,7 @@ class VirtualGPU : public device::VirtualDevice {
* used to synchronized on kernel outputs.
*/
bool hasPendingDispatch_;
amd::Monitor execution_; //!< Lock to serialise access to all device objects
Timestamp* timestamp_;
hsa_agent_t gpu_device_; //!< Physical device
hsa_queue_t* gpu_queue_; //!< Queue associated with a gpu
@@ -321,7 +327,7 @@ class VirtualGPU : public device::VirtualDevice {
uint kernarg_pool_cur_offset_;
std::vector<ProfilingSignal> signal_pool_; //!< Pool of signals for profiling
const uint index_; //!< Virtual gpu unique index
uint index_; //!< Virtual gpu unique index
friend class Timestamp;
// PM4 packet for gfx8 performance counter
rocclr/runtime/platform/command.cpp
+8 -3
View File
@@ -226,8 +226,13 @@ void Command::enqueue() {
const Context& Command::context() const { return queue_->context(); }
NDRangeKernelCommand::NDRangeKernelCommand(HostQueue& queue, const EventWaitList& eventWaitList,
Kernel& kernel, const NDRangeContainer& sizes, uint32_t sharedMemBytes)
: Command(queue, CL_COMMAND_NDRANGE_KERNEL, eventWaitList), kernel_(kernel), sizes_(sizes), sharedMemBytes_(sharedMemBytes) {
Kernel& kernel, const NDRangeContainer& sizes,
uint32_t sharedMemBytes, uint32_t extraParam)
: Command(queue, CL_COMMAND_NDRANGE_KERNEL, eventWaitList)
, kernel_(kernel)
, sizes_(sizes)
, sharedMemBytes_(sharedMemBytes)
, extraParam_(extraParam) {
auto& device = queue.device();
auto devKernel = const_cast<device::Kernel*>(kernel.getDeviceKernel(device));
profilingInfo_.setCallback(devKernel->getProfilingCallback(
@@ -393,7 +398,7 @@ bool MigrateMemObjectsCommand::validateMemory() {
cl_int NDRangeKernelCommand::captureAndValidate() {
const amd::Device& device = queue()->device();
// Validate the kernel before submission
if (!queue()->device().validateKernel(kernel(), queue()->vdev())) {
if (!queue()->device().validateKernel(kernel(), queue()->vdev(), cooperativeGroups())) {
return CL_OUT_OF_RESOURCES;
}
rocclr/runtime/platform/command.hpp
+15 -1
View File
@@ -753,11 +753,18 @@ class NDRangeKernelCommand : public Command {
NDRangeContainer sizes_;
address parameters_;
uint32_t sharedMemBytes_;
uint32_t extraParam_;
public:
enum {
CooperativeGroups = 0x01,
CooperativeMultiDeviceGroups = 0x02,
};
//! Construct an ExecuteKernel command
NDRangeKernelCommand(HostQueue& queue, const EventWaitList& eventWaitList, Kernel& kernel,
const NDRangeContainer& sizes, uint32_t sharedMemBytes = 0);
const NDRangeContainer& sizes, uint32_t sharedMemBytes = 0,
uint32_t extraParam = 0);
virtual void submit(device::VirtualDevice& device) { device.submitKernel(*this); }
@@ -776,6 +783,13 @@ class NDRangeKernelCommand : public Command {
//! Return the shared memory size
uint32_t sharedMemBytes() const { return sharedMemBytes_; }
//! Return the cooperative groups mode
bool cooperativeGroups() const { return (extraParam_ & CooperativeGroups) ? true : false; }
//! Return the cooperative multi device groups mode
bool cooperativeMultiDeviceGroups() const {
return (extraParam_ & CooperativeMultiDeviceGroups) ? true : false; }
//! Set the local work size.
void setLocalWorkSize(const NDRange& local) { sizes_.local() = local; }
rocclr/runtime/utils/flags.hpp
+2
View File
@@ -173,6 +173,8 @@ release(bool, GPU_ENABLE_LC, IS_LIGHTNING, \
"Enables LC path") \
release(bool, GPU_ENABLE_HW_P2P, false, \
"Enables HW P2P path") \
release(bool, GPU_ENABLE_COOP_GROUPS, false, \
"Enables cooperative group launch") \
release(uint, GPU_MAX_COMMAND_BUFFERS, 8, \
"The maximum number of command buffers allocated per queue") \
release(cstring, HIP_VISIBLE_DEVICES, "", \