diff --git a/projects/clr/rocclr/device/blit.cpp b/projects/clr/rocclr/device/blit.cpp index 7bbf73a192..15eeb89028 100644 --- a/projects/clr/rocclr/device/blit.cpp +++ b/projects/clr/rocclr/device/blit.cpp @@ -553,8 +553,8 @@ bool HostBlitManager::copyImage(device::Memory& srcMemory, device::Memory& dstMe } bool HostBlitManager::fillBuffer(device::Memory& memory, const void* pattern, size_t patternSize, - const amd::Coord3D& origin, const amd::Coord3D& size, - bool entire) const { + const amd::Coord3D& origin, const amd::Coord3D& size, bool entire, + bool forceBlit) const { // Map memory void* fillMem = memory.cpuMap(vDev_, (entire) ? Memory::CpuWriteOnly : 0); if (fillMem == NULL) { diff --git a/projects/clr/rocclr/device/blit.hpp b/projects/clr/rocclr/device/blit.hpp index 7865a036f9..b87f3d8542 100644 --- a/projects/clr/rocclr/device/blit.hpp +++ b/projects/clr/rocclr/device/blit.hpp @@ -179,9 +179,10 @@ class BlitManager : public amd::HeapObject { const void* pattern, //!< Pattern data size_t patternSize, //!< Pattern size const amd::Coord3D& origin, //!< Destination origin - const amd::Coord3D& size, //!< Size of the copy region - bool entire = false //!< Entire buffer will be updated - ) const = 0; + const amd::Coord3D& size, //!< Size of the fill region + bool entire = false, //!< Entire buffer will be updated + bool forceBlit = false //!< Force GPU Blit for fill + ) const = 0; //! Fills an image memory with a pattern data virtual bool fillImage(Memory& dstMemory, //!< Memory object to fill with pattern @@ -193,7 +194,7 @@ class BlitManager : public amd::HeapObject { //! Enables synchronization on blit operations void enableSynchronization() { syncOperation_ = true; } - + //! Returns Xfer queue lock virtual amd::Monitor* lockXfer() const { return nullptr; } @@ -331,9 +332,10 @@ class HostBlitManager : public device::BlitManager { const void* pattern, //!< Pattern data size_t patternSize, //!< Pattern size const amd::Coord3D& origin, //!< Destination origin - const amd::Coord3D& size, //!< Size of the copy region - bool entire = false //!< Entire buffer will be updated - ) const; + const amd::Coord3D& size, //!< Size of the fill region + bool entire = false, //!< Entire buffer will be updated + bool forceBlit = false //!< Force GPU Blit for fill + ) const; //! Fills an image memory with a pattern data virtual bool fillImage(device::Memory& dstMemory, //!< Memory object to fill with pattern diff --git a/projects/clr/rocclr/device/device.hpp b/projects/clr/rocclr/device/device.hpp index d5550dfe5a..4ac05cb2b0 100644 --- a/projects/clr/rocclr/device/device.hpp +++ b/projects/clr/rocclr/device/device.hpp @@ -80,6 +80,7 @@ class SvmMapMemoryCommand; class SvmUnmapMemoryCommand; class SvmPrefetchAsyncCommand; class TransferBufferFileCommand; +class StreamOperationCommand; class HwDebugManager; class Isa; class Device; @@ -1191,6 +1192,7 @@ class VirtualDevice : public amd::HeapObject { virtual void submitSvmPrefetchAsync(amd::SvmPrefetchAsyncCommand& cmd) { ShouldNotReachHere(); } + virtual void submitStreamOperation(amd::StreamOperationCommand& cmd) { ShouldNotReachHere(); } //! Get the blit manager object device::BlitManager& blitMgr() const { return *blitMgr_; } diff --git a/projects/clr/rocclr/device/gpu/gpublit.cpp b/projects/clr/rocclr/device/gpu/gpublit.cpp index 44c58cb01b..32cd4985d3 100644 --- a/projects/clr/rocclr/device/gpu/gpublit.cpp +++ b/projects/clr/rocclr/device/gpu/gpublit.cpp @@ -2024,12 +2024,12 @@ bool KernelBlitManager::writeBufferRect(const void* srcHost, device::Memory& dst bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern, size_t patternSize, const amd::Coord3D& origin, const amd::Coord3D& size, - bool entire) const { + bool entire, bool forceBlit) const { amd::ScopedLock k(lockXferOps_); bool result = false; // Use host fill if memory has direct access - if (setup_.disableFillBuffer_ || gpuMem(memory).isHostMemDirectAccess()) { + if (setup_.disableFillBuffer_ || (!forceBlit && gpuMem(memory).isHostMemDirectAccess())) { result = HostBlitManager::fillBuffer(memory, pattern, patternSize, origin, size, entire); synchronize(); return result; diff --git a/projects/clr/rocclr/device/gpu/gpublit.hpp b/projects/clr/rocclr/device/gpu/gpublit.hpp index 6f3135160f..913f850978 100644 --- a/projects/clr/rocclr/device/gpu/gpublit.hpp +++ b/projects/clr/rocclr/device/gpu/gpublit.hpp @@ -345,9 +345,10 @@ class KernelBlitManager : public DmaBlitManager { const void* pattern, //!< Pattern data size_t patternSize, //!< Pattern size const amd::Coord3D& origin, //!< Destination origin - const amd::Coord3D& size, //!< Size of the copy region - bool entire = false //!< Entire buffer will be updated - ) const; + const amd::Coord3D& size, //!< Size of the fill region + bool entire = false, //!< Entire buffer will be updated + bool forceBlit = false //!< Force GPU Blit for fill + ) const; //! Fills an image memory with a pattern data virtual bool fillImage(device::Memory& dstMemory, //!< Memory object to fill with pattern diff --git a/projects/clr/rocclr/device/pal/palblit.cpp b/projects/clr/rocclr/device/pal/palblit.cpp index 23291e209c..55cac9adec 100644 --- a/projects/clr/rocclr/device/pal/palblit.cpp +++ b/projects/clr/rocclr/device/pal/palblit.cpp @@ -2065,12 +2065,12 @@ bool KernelBlitManager::writeBufferRect(const void* srcHost, device::Memory& dst bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern, size_t patternSize, const amd::Coord3D& origin, const amd::Coord3D& size, - bool entire) const { + bool entire, bool forceBlit) const { amd::ScopedLock k(lockXferOps_); bool result = false; // Use host fill if memory has direct access - if (setup_.disableFillBuffer_ || gpuMem(memory).isHostMemDirectAccess()) { + if (setup_.disableFillBuffer_ || (!forceBlit && gpuMem(memory).isHostMemDirectAccess())) { result = HostBlitManager::fillBuffer(memory, pattern, patternSize, origin, size, entire); synchronize(); return result; diff --git a/projects/clr/rocclr/device/pal/palblit.hpp b/projects/clr/rocclr/device/pal/palblit.hpp index 9f1198c0c8..571ce80207 100644 --- a/projects/clr/rocclr/device/pal/palblit.hpp +++ b/projects/clr/rocclr/device/pal/palblit.hpp @@ -345,9 +345,10 @@ class KernelBlitManager : public DmaBlitManager { const void* pattern, //!< Pattern data size_t patternSize, //!< Pattern size const amd::Coord3D& origin, //!< Destination origin - const amd::Coord3D& size, //!< Size of the copy region - bool entire = false //!< Entire buffer will be updated - ) const; + const amd::Coord3D& size, //!< Size of the fill region + bool entire = false, //!< Entire buffer will be updated + bool forceBlit = false //!< Force GPU Blit for fill + ) const; //! Fills an image memory with a pattern data virtual bool fillImage(device::Memory& dstMemory, //!< Memory object to fill with pattern diff --git a/projects/clr/rocclr/device/rocm/rocblit.cpp b/projects/clr/rocclr/device/rocm/rocblit.cpp index c010500c0c..39c8feab29 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.cpp +++ b/projects/clr/rocclr/device/rocm/rocblit.cpp @@ -1949,12 +1949,12 @@ bool KernelBlitManager::writeBufferRect(const void* srcHost, device::Memory& dst // ================================================================================================ bool KernelBlitManager::fillBuffer(device::Memory& memory, const void* pattern, size_t patternSize, const amd::Coord3D& origin, const amd::Coord3D& size, - bool entire) const { + bool entire, bool forceBlit) const { amd::ScopedLock k(lockXferOps_); bool result = false; // Use host fill if memory has direct access - if (setup_.disableFillBuffer_ || memory.isHostMemDirectAccess()) { + if (setup_.disableFillBuffer_ || (!forceBlit && memory.isHostMemDirectAccess())) { // Stall GPU before CPU access gpu().releaseGpuMemoryFence(); result = HostBlitManager::fillBuffer(memory, pattern, patternSize, origin, size, entire); diff --git a/projects/clr/rocclr/device/rocm/rocblit.hpp b/projects/clr/rocclr/device/rocm/rocblit.hpp index a81480f21d..24959350ec 100644 --- a/projects/clr/rocclr/device/rocm/rocblit.hpp +++ b/projects/clr/rocclr/device/rocm/rocblit.hpp @@ -364,9 +364,10 @@ class KernelBlitManager : public DmaBlitManager { const void* pattern, //!< Pattern data size_t patternSize, //!< Pattern size const amd::Coord3D& origin, //!< Destination origin - const amd::Coord3D& size, //!< Size of the copy region - bool entire = false //!< Entire buffer will be updated - ) const; + const amd::Coord3D& size, //!< Size of the fill region + bool entire = false, //!< Entire buffer will be updated + bool forceBlit = false //!< Force GPU Blit for fill + ) const; //! Fills an image memory with a pattern data virtual bool fillImage(device::Memory& dstMemory, //!< Memory object to fill with pattern diff --git a/projects/clr/rocclr/device/rocm/rocdevice.cpp b/projects/clr/rocclr/device/rocm/rocdevice.cpp index a0ed5b46f0..89c118c225 100644 --- a/projects/clr/rocclr/device/rocm/rocdevice.cpp +++ b/projects/clr/rocclr/device/rocm/rocdevice.cpp @@ -2414,8 +2414,10 @@ bool Device::SetClockMode(const cl_set_device_clock_mode_input_amd setClockModeI static void callbackQueue(hsa_status_t status, hsa_queue_t* queue, void* data) { if (status != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK) { // Abort on device exceptions. - ClPrint(amd::LOG_NONE, amd::LOG_ALWAYS, "Device::callbackQueue aborting with status: 0x%x", - status); + const char* errorMsg = 0; + hsa_status_string(status, &errorMsg); + ClPrint(amd::LOG_NONE, amd::LOG_ALWAYS, + "Device::callbackQueue aborting with error : %s code: 0x%x", errorMsg, status); abort(); } } diff --git a/projects/clr/rocclr/device/rocm/rocmemory.cpp b/projects/clr/rocclr/device/rocm/rocmemory.cpp index 4aeb47c42f..2f7c74dcc6 100644 --- a/projects/clr/rocclr/device/rocm/rocmemory.cpp +++ b/projects/clr/rocclr/device/rocm/rocmemory.cpp @@ -639,6 +639,12 @@ void Buffer::destroy() { #else dev().hostFree(deviceMemory_, size());; #endif // AMD_HMM_SUPPORT + } else if (memFlags & ROCCLR_MEM_HSA_SIGNAL_MEMORY) { + if (HSA_STATUS_SUCCESS != hsa_signal_destroy(signal_)) { + ClPrint(amd::LOG_DEBUG, amd::LOG_MEM, + "[ROCClr] ROCCLR_MEM_HSA_SIGNAL_MEMORY signal destroy failed \n"); + } + deviceMemory_ = nullptr; } else { dev().hostFree(deviceMemory_, size()); } @@ -734,6 +740,28 @@ bool Buffer::create() { #endif // AMD_HMM_SUPPORT } else if (memFlags & CL_MEM_FOLLOW_USER_NUMA_POLICY) { deviceMemory_ = dev().hostNumaAlloc(size(), 1, (memFlags & CL_MEM_SVM_ATOMICS) != 0); + } else if (memFlags & ROCCLR_MEM_HSA_SIGNAL_MEMORY) { + // TODO: ROCr will introduce a new attribute enum that implies a non-blocking signal, + // replace "HSA_AMD_SIGNAL_AMD_GPU_ONLY" with this new enum when it is ready. + if (HSA_STATUS_SUCCESS != + hsa_amd_signal_create(kInitSignalValueOne, 0, nullptr, HSA_AMD_SIGNAL_AMD_GPU_ONLY, + &signal_)) { + ClPrint(amd::LOG_ERROR, amd::LOG_MEM, + "[ROCclr] ROCCLR_MEM_HSA_SIGNAL_MEMORY signal creation failed"); + return false; + } + volatile hsa_signal_value_t* signalValuePtr; + if (HSA_STATUS_SUCCESS != hsa_amd_signal_value_pointer(signal_, &signalValuePtr)) { + ClPrint(amd::LOG_ERROR, amd::LOG_MEM, + "[ROCclr] ROCCLR_MEM_HSA_SIGNAL_MEMORY pointer query failed"); + return false; + } + + deviceMemory_ = const_cast(signalValuePtr); // conversion to void * is + // implicit + + // Disable host access to force blit path for memeory writes. + flags_ &= ~HostMemoryDirectAccess; } else { deviceMemory_ = dev().hostAlloc(size(), 1, ((memFlags & CL_MEM_SVM_ATOMICS) != 0) ? Device::MemorySegment::kAtomics diff --git a/projects/clr/rocclr/device/rocm/rocmemory.hpp b/projects/clr/rocclr/device/rocm/rocmemory.hpp index b341c379e1..1bb5a1f52a 100644 --- a/projects/clr/rocclr/device/rocm/rocmemory.hpp +++ b/projects/clr/rocclr/device/rocm/rocmemory.hpp @@ -165,7 +165,13 @@ class Buffer : public roc::Memory { // Recreate the device memory using new size and alignment. bool recreate(size_t newSize, size_t newAlignment, bool forceSystem); + // Returns the HSA signal associated with this Memory object. + hsa_signal_t getSignal() const { return signal_; } + private: + // signal object used when ROCCLR_MEM_HSA_SIGNAL_MEMORY is set + hsa_signal_t signal_; + // Disable copy constructor Buffer(const Buffer&); diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp index e61ef4834d..373a597113 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp @@ -27,6 +27,7 @@ #include "platform/kernel.hpp" #include "platform/context.hpp" #include "platform/command.hpp" +#include "platform/command_utils.hpp" #include "platform/memory.hpp" #include "platform/sampler.hpp" #include "rochostcall.hpp" @@ -90,6 +91,11 @@ static constexpr uint16_t kBarrierPacketReleaseHeader = (HSA_FENCE_SCOPE_NONE << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) | (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE); +static constexpr uint16_t kBarrierVendorPacketHeader = + (HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE) | (1 << HSA_PACKET_HEADER_BARRIER) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE); + static constexpr hsa_barrier_and_packet_t kBarrierAcquirePacket = { kBarrierPacketAcquireHeader, 0, 0, {{0}}, 0, {0}}; @@ -2069,6 +2075,110 @@ void VirtualGPU::submitFillMemory(amd::FillMemoryCommand& cmd) { profilingEnd(cmd); } +void VirtualGPU::dispatchBarrierValuePacket(const hsa_amd_barrier_value_packet_t* packet, + hsa_amd_vendor_packet_header_t header) { + assert(packet->completion_signal.handle != 0); + const uint32_t queueSize = gpu_queue_->size; + const uint32_t queueMask = queueSize - 1; + + uint64_t index = hsa_queue_add_write_index_screlease(gpu_queue_, 1); + while ((index - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= queueMask) { + amd::Os::yield(); + } + hsa_amd_barrier_value_packet_t* aql_loc = &(reinterpret_cast( + gpu_queue_->base_address))[index & queueMask]; + *aql_loc = *packet; + unsigned int* headerPtr = reinterpret_cast(&header); + __atomic_store_n(reinterpret_cast(aql_loc), *headerPtr, __ATOMIC_RELEASE); + + hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index); + ClPrint(amd::LOG_DEBUG, amd::LOG_AQL, + "[%zx] HWq=0x%zx, BarrierValue Header = 0x%x AmdFormat = 0x%x ", + "(type=%d, barrier=%d, acquire=%d, release=%d), " + "completion_signal=0x%zx value = 0x%llx mask = 0x%llx cond: %d (GTE: %d EQ: %d NE: %d)", + std::this_thread::get_id(), gpu_queue_, header.header, header.AmdFormat, + extractAqlBits(header.header, HSA_PACKET_HEADER_TYPE, HSA_PACKET_HEADER_WIDTH_TYPE), + extractAqlBits(header.header, HSA_PACKET_HEADER_BARRIER, HSA_PACKET_HEADER_WIDTH_BARRIER), + extractAqlBits(header.header, HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE, + HSA_PACKET_HEADER_WIDTH_SCACQUIRE_FENCE_SCOPE), + extractAqlBits(header.header, HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE, + HSA_PACKET_HEADER_WIDTH_SCRELEASE_FENCE_SCOPE), + packet->completion_signal, packet->value, packet->mask, packet->cond, + HSA_SIGNAL_CONDITION_GTE, HSA_SIGNAL_CONDITION_EQ, HSA_SIGNAL_CONDITION_NE); +} + +void VirtualGPU::submitStreamOperation(amd::StreamOperationCommand& cmd) { + // Make sure VirtualGPU has an exclusive access to the resources + amd::ScopedLock lock(execution()); + profilingBegin(cmd); + + const cl_command_type type = cmd.type(); + const int64_t value = cmd.value(); + const uint64_t mask = cmd.mask(); + const unsigned int flags = cmd.flags(); + const size_t sizeBytes = cmd.sizeBytes(); + const size_t offset = cmd.offset(); + + amd::Memory* amdMemory = &cmd.memory(); + Memory* memory = dev().getRocMemory(amdMemory); + + if (type == ROCCLR_COMMAND_STREAM_WAIT_VALUE) { + hsa_amd_barrier_value_packet_t aqlPacket; + hsa_amd_vendor_packet_header_t header; + hsa_signal_t signal; + Buffer* buff = static_cast(memory); + + header.header = kBarrierVendorPacketHeader; + header.AmdFormat = HSA_AMD_PACKET_TYPE_BARRIER_VALUE; + aqlPacket.signal = buff->getSignal(); + aqlPacket.completion_signal = Barriers().ActiveSignal(); + + // mask is always applied on value at signal before performing + // the comparision defiend by 'condition' + switch (flags) { + case ROCCLR_STREAM_WAIT_VALUE_GTE: + aqlPacket.value = value; + aqlPacket.mask = mask; + aqlPacket.cond = HSA_SIGNAL_CONDITION_GTE; + break; + case ROCCLR_STREAM_WAIT_VALUE_EQ: + aqlPacket.value = value; + aqlPacket.mask = mask; + aqlPacket.cond = HSA_SIGNAL_CONDITION_EQ; + break; + case ROCCLR_STREAM_WAIT_VALUE_AND: + aqlPacket.value = 0; + aqlPacket.mask = (value & mask); + aqlPacket.cond = HSA_SIGNAL_CONDITION_NE; + break; + case ROCCLR_STREAM_WAIT_VALUE_NOR: + aqlPacket.value = ~value & mask; + aqlPacket.mask = ~value & mask; + aqlPacket.cond = HSA_SIGNAL_CONDITION_NE; + break; + default: + ShouldNotReachHere(); + break; + } + dispatchBarrierValuePacket(&aqlPacket, header); + } else if (type == ROCCLR_COMMAND_STREAM_WRITE_VALUE) { + amd::Coord3D origin(offset); + amd::Coord3D size(sizeBytes); + bool entire = amdMemory->isEntirelyCovered(origin, size); + + // Use GPU Blit to write + bool result = blitMgr().fillBuffer(*memory, &value, sizeBytes, origin, size, entire, true); + ClPrint(amd::LOG_DEBUG, amd::LOG_COPY, "Writting value: 0x%lx \n", value); + + if (!result) { + LogError("submitStreamOperation: Write failed!"); + } + } else { + ShouldNotReachHere(); + } + profilingEnd(cmd); +} + void VirtualGPU::submitSvmFillMemory(amd::SvmFillMemoryCommand& cmd) { // Make sure VirtualGPU has an exclusive access to the resources amd::ScopedLock lock(execution()); @@ -2900,7 +3010,6 @@ void VirtualGPU::submitPerfCounter(amd::PerfCounterCommand& vcmd) { // Make sure all performance counter objects to use the same profile PerfCounter* counter = nullptr; for (uint i = 0; i < vcmd.getNumCounters(); ++i) { - amd::PerfCounter* amdCounter = static_cast(counters[i]); counter = static_cast(amdCounter->getDeviceCounter()); diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.hpp b/projects/clr/rocclr/device/rocm/rocvirtual.hpp index 626a684efd..2ebf5ede6d 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.hpp +++ b/projects/clr/rocclr/device/rocm/rocvirtual.hpp @@ -251,6 +251,7 @@ class VirtualGPU : public device::VirtualDevice { void flush(amd::Command* list = nullptr, bool wait = false); void submitFillMemory(amd::FillMemoryCommand& cmd); + void submitStreamOperation(amd::StreamOperationCommand& cmd); void submitMigrateMemObjects(amd::MigrateMemObjectsCommand& cmd); void submitSvmFreeMemory(amd::SvmFreeMemoryCommand& cmd); @@ -335,6 +336,8 @@ class VirtualGPU : public device::VirtualDevice { bool skipSignal = false); bool dispatchCounterAqlPacket(hsa_ext_amd_aql_pm4_packet_t* packet, const uint32_t gfxVersion, bool blocking, const hsa_ven_amd_aqlprofile_1_00_pfn_t* extApi); + void dispatchBarrierValuePacket(const hsa_amd_barrier_value_packet_t* packet, + hsa_amd_vendor_packet_header_t header); void initializeDispatchPacket(hsa_kernel_dispatch_packet_t* packet, amd::NDRangeContainer& sizes); diff --git a/projects/clr/rocclr/platform/command.hpp b/projects/clr/rocclr/platform/command.hpp index 8e6a78ce96..b721583517 100644 --- a/projects/clr/rocclr/platform/command.hpp +++ b/projects/clr/rocclr/platform/command.hpp @@ -42,6 +42,7 @@ #include "platform/perfctr.hpp" #include "platform/threadtrace.hpp" #include "platform/activity.hpp" +#include "platform/command_utils.hpp" #include "CL/cl_ext.h" @@ -599,6 +600,58 @@ class FillMemoryCommand : public OneMemoryArgCommand { bool isEntireMemory() const; }; +/*! \brief A stream operation command. + * + * \details Used to perform a stream wait or strem write operations. + * Wait: All the commands issued after stream wait are not executed until the wait + * condition is true. + * Write: Writes a 32 or 64 bit vaue to the memeory using a GPU Blit. + */ + +class StreamOperationCommand : public OneMemoryArgCommand { + private: + int64_t value_; // !< Value to Wait on or to Write. + uint64_t mask_; // !< Mask to be applied on signal value for Wait operation. + unsigned int flags_; // !< Flags defining the Wait condition. + size_t offset_; // !< Offset into memory for Write + size_t sizeBytes_; // !< Size in bytes to Write. + + // NOTE: mask_ is only used for wait operation and + // offset and sizeBytes are only used for write. + + public: + StreamOperationCommand(HostQueue& queue, cl_command_type cmdType, + const EventWaitList& eventWaitList, Memory& memory, const int64_t value, + const uint64_t mask, unsigned int flags, size_t offset, size_t sizeBytes) + : OneMemoryArgCommand(queue, cmdType, eventWaitList, memory), + value_(value), + mask_(mask), + flags_(flags), + offset_(offset), + sizeBytes_(sizeBytes) { + // Sanity check + assert((cmdType == ROCCLR_COMMAND_STREAM_WRITE_VALUE || + (cmdType == ROCCLR_COMMAND_STREAM_WAIT_VALUE && + memory_->getMemFlags() & ROCCLR_MEM_HSA_SIGNAL_MEMORY)) && + "Invalid Stream Operation"); + } + + virtual void submit(device::VirtualDevice& device) { device.submitStreamOperation(*this); } + + //! Returns the value + const int64_t value() const { return value_; } + //! Returns the wait mask + const uint64_t mask() const { return mask_; } + //! Return the wait flags + const unsigned int flags() const { return flags_; } + //! Return the memory object. + Memory& memory() const { return *memory_; } + //! Return the write offset. + const size_t offset() const { return offset_; } + //! Return the write size. + const size_t sizeBytes() const { return sizeBytes_; } +}; + /*! \brief A generic copy memory command * * \details Used for both buffers and images. Backends are expected diff --git a/projects/clr/rocclr/platform/command_utils.hpp b/projects/clr/rocclr/platform/command_utils.hpp new file mode 100644 index 0000000000..c6b5304eca --- /dev/null +++ b/projects/clr/rocclr/platform/command_utils.hpp @@ -0,0 +1,9 @@ +// Dummy command types for Stream Wait and Write commands. +#define ROCCLR_COMMAND_STREAM_WAIT_VALUE 0x4501 +#define ROCCLR_COMMAND_STREAM_WRITE_VALUE 0x4502 + +// Stream Wait Value Conidtions +#define ROCCLR_STREAM_WAIT_VALUE_GTE 0x0 +#define ROCCLR_STREAM_WAIT_VALUE_EQ 0x1 +#define ROCCLR_STREAM_WAIT_VALUE_AND 0x2 +#define ROCCLR_STREAM_WAIT_VALUE_NOR 0x3 diff --git a/projects/clr/rocclr/platform/memory.hpp b/projects/clr/rocclr/platform/memory.hpp index c134056365..7c7966dcb5 100644 --- a/projects/clr/rocclr/platform/memory.hpp +++ b/projects/clr/rocclr/platform/memory.hpp @@ -37,6 +37,7 @@ #include #include #define CL_MEM_FOLLOW_USER_NUMA_POLICY (1u << 31) +#define ROCCLR_MEM_HSA_SIGNAL_MEMORY (1u << 30) namespace device { class Memory;