AI/rocm-systems
2
0
Atdalīts 0
Kods Problēmas Izmaiņu pieprasījumi Darbības Pakotnes Projekti Laidieni Vikivietne Aktivitāte

SWDEV-179105 - Stream Operations: Add support for Wait and Write

Pārlūkot izejas kodu 
Change-Id: Ibffa1d6d573826b64763da280074a77271d66808


[ROCm/clr commit: 0a5f9a3b10]
Šī revīzija ir iekļauta:
Ravi C Akkenapally
2021-02-08 21:51:18 -08:00
vecāks 4583cbafee
revīzija 6629930067
17 mainīti faili ar 245 papildinājumiem un 27 dzēšanām
Lejupielādēt ielāpa failu Lejupielādēt izmaiņu failu Izvērst visus failus Savērst visus failus
projects/clr/rocclr/device/blit.cpp
+2 -2
Parādīt failu
@@ -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) {
projects/clr/rocclr/device/blit.hpp
+9 -7
Parādīt failu
@@ -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
projects/clr/rocclr/device/device.hpp
+2
Parādīt failu
@@ -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_; }
projects/clr/rocclr/device/gpu/gpublit.cpp
+2 -2
Parādīt failu
@@ -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;
projects/clr/rocclr/device/gpu/gpublit.hpp
+4 -3
Parādīt failu
@@ -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
projects/clr/rocclr/device/pal/palblit.cpp
+2 -2
Parādīt failu
@@ -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;
projects/clr/rocclr/device/pal/palblit.hpp
+4 -3
Parādīt failu
@@ -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
projects/clr/rocclr/device/rocm/rocblit.cpp
+2 -2
Parādīt failu
@@ -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);
projects/clr/rocclr/device/rocm/rocblit.hpp
+4 -3
Parādīt failu
@@ -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
projects/clr/rocclr/device/rocm/rocdevice.cpp
+4 -2
Parādīt failu
@@ -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();
}
}
projects/clr/rocclr/device/rocm/rocmemory.cpp
+28
Parādīt failu
@@ -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<long int*>(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
projects/clr/rocclr/device/rocm/rocmemory.hpp
+6
Parādīt failu
@@ -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&);
projects/clr/rocclr/device/rocm/rocvirtual.cpp
+110 -1
Parādīt failu
@@ -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<hsa_amd_barrier_value_packet_t*>(
gpu_queue_->base_address))[index & queueMask];
*aql_loc = *packet;
unsigned int* headerPtr = reinterpret_cast<unsigned int*>(&header);
__atomic_store_n(reinterpret_cast<uint32_t*>(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<Buffer*>(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<amd::PerfCounter*>(counters[i]);
counter = static_cast<PerfCounter*>(amdCounter->getDeviceCounter());
projects/clr/rocclr/device/rocm/rocvirtual.hpp
+3
Parādīt failu
@@ -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);
projects/clr/rocclr/platform/command.hpp
+53
Parādīt failu
@@ -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
projects/clr/rocclr/platform/command_utils.hpp
+9
Parādīt failu
@@ -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
projects/clr/rocclr/platform/memory.hpp
+1
Parādīt failu
@@ -37,6 +37,7 @@
#include <unordered_map>
#include <memory>
#define CL_MEM_FOLLOW_USER_NUMA_POLICY (1u << 31)
#define ROCCLR_MEM_HSA_SIGNAL_MEMORY (1u << 30)
namespace device {
class Memory;