Check open source core runtime code into perforce. This includes license and README files.

[git-p4: depot-paths = "//depot/stg/hsa/drivers/hsa/runtime/": change = 1249136]
This commit is contained in:
James Edwards (xN/A) TX
2016-03-20 15:39:40 -05:00
parent 73d43224e9
commit c9ffe0004e
100 changed files with 37877 additions and 0 deletions
+264
View File
@@ -0,0 +1,264 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
// HSA runtime C++ interface file.
#ifndef HSA_RUNTME_CORE_INC_AGENT_H_
#define HSA_RUNTME_CORE_INC_AGENT_H_
#include <assert.h>
#include <vector>
#include "core/inc/runtime.h"
#include "core/inc/checked.h"
#include "core/inc/isa.h"
#include "core/inc/queue.h"
#include "core/inc/memory_region.h"
#include "core/util/utils.h"
namespace core {
class Signal;
typedef void (*HsaEventCallback)(hsa_status_t status, hsa_queue_t* source,
void* data);
class MemoryRegion;
// Agent is intended to be an pure interface class and may be wrapped or
// replaced by tools libraries. All funtions other than Convert, node_id,
// device_type, and public_handle must be virtual.
class Agent : public Checked<0xF6BC25EB17E6F917> {
public:
// @brief Convert agent object into hsa_agent_t.
//
// @param [in] agent Pointer to an agent.
//
// @retval hsa_agent_t
static __forceinline hsa_agent_t Convert(Agent* agent) {
const hsa_agent_t agent_handle = {
static_cast<uint64_t>(reinterpret_cast<uintptr_t>(agent))};
return agent_handle;
}
// @brief Convert agent object into const hsa_agent_t.
//
// @param [in] agent Pointer to an agent.
//
// @retval const hsa_agent_t
static __forceinline const hsa_agent_t Convert(const Agent* agent) {
const hsa_agent_t agent_handle = {
static_cast<uint64_t>(reinterpret_cast<uintptr_t>(agent))};
return agent_handle;
}
// @brief Convert hsa_agent_t handle into Agent*.
//
// @param [in] agent An hsa_agent_t handle.
//
// @retval Agent*
static __forceinline Agent* Convert(hsa_agent_t agent) {
return reinterpret_cast<Agent*>(agent.handle);
}
// Lightweight RTTI for vendor specific implementations.
enum DeviceType { kAmdGpuDevice = 0, kAmdCpuDevice = 1, kUnknownDevice = 2 };
// @brief Agent class contructor.
//
// @param [in] type CPU or GPU or other.
explicit Agent(uint32_t node_id, DeviceType type)
: node_id_(node_id), device_type_(uint32_t(type)) {
public_handle_ = Convert(this);
}
// @brief Agent class contructor.
//
// @param [in] type CPU or GPU or other.
explicit Agent(uint32_t node_id, uint32_t type)
: node_id_(node_id), device_type_(type) {
public_handle_ = Convert(this);
}
// @brief Agent class destructor.
virtual ~Agent() {}
// @brief Submit DMA copy command to move data from src to dst and wait
// until it is finished.
//
// @details The agent must be able to access @p dst and @p src.
//
// @param [in] dst Memory address of the destination.
// @param [in] src Memory address of the source.
// @param [in] size Copy size in bytes.
//
// @retval HSA_STATUS_SUCCESS The memory copy is finished and successful.
virtual hsa_status_t DmaCopy(void* dst, const void* src, size_t size) {
return HSA_STATUS_ERROR;
}
// @brief Submit DMA copy command to move data from src to dst. This call
// does not wait until the copy is finished
//
// @details The agent must be able to access @p dst and @p src. Memory copy
// will be performed after all signals in @p dep_signals have value of 0.
// On memory copy completion, the value of out_signal is decremented.
//
// @param [in] dst Memory address of the destination.
// @param [in] src Memory address of the source.
// @param [in] size Copy size in bytes.
// @param [in] dep_signals Array of signal dependency.
// @param [in] out_signal Completion signal.
//
// @retval HSA_STATUS_SUCCESS The memory copy is finished and successful.
virtual hsa_status_t DmaCopy(void* dst, const void* src, size_t size,
std::vector<core::Signal*>& dep_signals,
core::Signal& out_signal) {
return HSA_STATUS_ERROR;
}
// @brief Submit DMA command to set the content of a pointer and wait
// until it is finished.
//
// @details The agent must be able to access @p ptr
//
// @param [in] ptr Address of the memory to be set.
// @param [in] value The value/pattern that will be used to set @p ptr.
// @param [in] count Number of uint32_t element to be set.
//
// @retval HSA_STATUS_SUCCESS The memory fill is finished and successful.
virtual hsa_status_t DmaFill(void* ptr, uint32_t value, size_t count) {
return HSA_STATUS_ERROR;
}
// @brief Invoke the user provided callback for each region accessible by
// this agent.
//
// @param [in] callback User provided callback function.
// @param [in] data User provided pointer as input for @p callback.
//
// @retval ::HSA_STATUS_SUCCESS if the callback function for each traversed
// region returns ::HSA_STATUS_SUCCESS.
virtual hsa_status_t IterateRegion(
hsa_status_t (*callback)(hsa_region_t region, void* data),
void* data) const = 0;
// @brief Create queue.
//
// @param [in] size Number of packets the queue is expected to hold. Must be a
// power of 2 greater than 0.
// @param [in] queue_type Queue type.
// @param [in] event_callback Callback invoked for every
// asynchronous event related to the newly created queue. May be NULL.The HSA
// runtime passes three arguments to the callback : a code identifying the
// event that triggered the invocation, a pointer to the queue where the event
// originated, and the application data.
// @param [in] data Application data that is passed to @p callback.
// @param [in] private_segment_size A hint to indicate the maximum expected
// private segment usage per work-item, in bytes.
// @param [in] group_segment_size A hint to indicate the maximum expected
// group segment usage per work-group, in bytes.
// @param[out] queue Memory location where the HSA runtime stores a pointer
// to the newly created queue.
//
// @retval HSA_STATUS_SUCCESS The queue has been created successfully.
virtual hsa_status_t QueueCreate(size_t size, hsa_queue_type_t queue_type,
HsaEventCallback event_callback, void* data,
uint32_t private_segment_size,
uint32_t group_segment_size,
Queue** queue) = 0;
// @brief Query the value of an attribute.
//
// @param [in] attribute Attribute to query.
// @param [out] value Pointer to store the value of the attribute.
//
// @param HSA_STATUS_SUCCESS @p value has been filled with the value of the
// attribute.
virtual hsa_status_t GetInfo(hsa_agent_info_t attribute,
void* value) const = 0;
// @brief Returns an array of regions owned by the agent.
virtual const std::vector<const core::MemoryRegion*>& regions() const = 0;
// @details Returns the agent's instruction set architecture.
virtual const Isa* isa() const = 0;
// @brief Returns the device type (CPU/GPU/Others).
__forceinline uint32_t device_type() const { return device_type_; }
// @brief Returns hsa_agent_t handle exposed to end user.
//
// @details Only matters when tools library need to intercept HSA calls.
__forceinline hsa_agent_t public_handle() const { return public_handle_; }
// @brief Returns node id associated with this agent.
__forceinline uint32_t node_id() const { return node_id_; }
protected:
// Intention here is to have a polymorphic update procedure for public_handle_
// which is callable on any Agent* but only from some class dervied from
// Agent*. do_set_public_handle should remain protected or private in all
// derived types.
static __forceinline void set_public_handle(Agent* agent,
hsa_agent_t handle) {
agent->do_set_public_handle(handle);
}
virtual void do_set_public_handle(hsa_agent_t handle) {
public_handle_ = handle;
}
hsa_agent_t public_handle_;
private:
// @brief Node id.
const uint32_t node_id_;
const uint32_t device_type_;
// Forbid copying and moving of this object
DISALLOW_COPY_AND_ASSIGN(Agent);
};
} // namespace core
#endif // header guard
@@ -0,0 +1,412 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef HSA_RUNTIME_CORE_INC_AMD_HW_AQL_COMMAND_PROCESSOR_H_
#define HSA_RUNTIME_CORE_INC_AMD_HW_AQL_COMMAND_PROCESSOR_H_
#include "core/inc/runtime.h"
#include "core/inc/signal.h"
#include "core/inc/queue.h"
#include "core/inc/amd_gpu_agent.h"
namespace amd {
/// @brief Encapsulates HW Aql Command Processor functionality. It
/// provide the interface for things such as Doorbell register, read,
/// write pointers and a buffer.
class AqlQueue : public core::Queue, public core::Signal {
public:
static __forceinline bool IsType(core::Signal* signal) {
return signal->IsType(&rtti_id_);
}
// Acquires/releases queue resources and requests HW schedule/deschedule.
AqlQueue(GpuAgent* agent, size_t req_size_pkts, HSAuint32 node_id,
ScratchInfo& scratch, core::HsaEventCallback callback,
void* err_data, bool is_kv = false);
~AqlQueue();
/// @brief Indicates if queue is valid or not
bool IsValid() const { return valid_; }
/// @brief Queue interfaces
hsa_status_t Inactivate();
/// @brief Atomically reads the Read index of with Acquire semantics
///
/// @return uint64_t Value of read index
uint64_t LoadReadIndexAcquire();
/// @brief Atomically reads the Read index of with Relaxed semantics
///
/// @return uint64_t Value of read index
uint64_t LoadReadIndexRelaxed();
/// @brief Atomically reads the Write index of with Acquire semantics
///
/// @return uint64_t Value of write index
uint64_t LoadWriteIndexAcquire();
/// @brief Atomically reads the Write index of with Relaxed semantics
///
/// @return uint64_t Value of write index
uint64_t LoadWriteIndexRelaxed();
/// @brief This operation is illegal
void StoreReadIndexRelaxed(uint64_t value) { assert(false); }
/// @brief This operation is illegal
void StoreReadIndexRelease(uint64_t value) { assert(false); }
/// @brief Atomically writes the Write index of with Relaxed semantics
///
/// @param value New value of write index to update with
void StoreWriteIndexRelaxed(uint64_t value);
/// @brief Atomically writes the Write index of with Release semantics
///
/// @param value New value of write index to update with
void StoreWriteIndexRelease(uint64_t value);
/// @brief Compares and swaps Write index using Acquire and Release semantics
///
/// @param expected Current value of write index
///
/// @param value Value of new write index
///
/// @return uint64_t Value of write index before the update
uint64_t CasWriteIndexAcqRel(uint64_t expected, uint64_t value);
/// @brief Compares and swaps Write index using Acquire semantics
///
/// @param expected Current value of write index
///
/// @param value Value of new write index
///
/// @return uint64_t Value of write index before the update
uint64_t CasWriteIndexAcquire(uint64_t expected, uint64_t value);
/// @brief Compares and swaps Write index using Relaxed semantics
///
/// @param expected Current value of write index
///
/// @param value Value of new write index
///
/// @return uint64_t Value of write index before the update
uint64_t CasWriteIndexRelaxed(uint64_t expected, uint64_t value);
/// @brief Compares and swaps Write index using Release semantics
///
/// @param expected Current value of write index
///
/// @param value Value of new write index
///
/// @return uint64_t Value of write index before the update
uint64_t CasWriteIndexRelease(uint64_t expected, uint64_t value);
/// @brief Updates the Write index using Acquire and Release semantics
///
/// @param value Value of new write index
///
/// @return uint64_t Value of write index before the update
uint64_t AddWriteIndexAcqRel(uint64_t value);
/// @brief Updates the Write index using Acquire semantics
///
/// @param value Value of new write index
///
/// @return uint64_t Value of write index before the update
uint64_t AddWriteIndexAcquire(uint64_t value);
/// @brief Updates the Write index using Relaxed semantics
///
/// @param value Value of new write index
///
/// @return uint64_t Value of write index before the update
uint64_t AddWriteIndexRelaxed(uint64_t value);
/// @brief Updates the Write index using Release semantics
///
/// @param value Value of new write index
///
/// @return uint64_t Value of write index before the update
uint64_t AddWriteIndexRelease(uint64_t value);
/// @brief Set CU Masking
///
/// @param num_cu_mask_count size of mask bit array
///
/// @param cu_mask pointer to cu mask
///
/// @return hsa_status_t
hsa_status_t SetCUMasking(const uint32_t num_cu_mask_count,
const uint32_t* cu_mask);
/// @brief This operation is illegal
hsa_signal_value_t LoadRelaxed() {
assert(false);
return 0;
}
/// @brief This operation is illegal
hsa_signal_value_t LoadAcquire() {
assert(false);
return 0;
}
/// @brief Update signal value using Relaxed semantics
void StoreRelaxed(hsa_signal_value_t value);
/// @brief Update signal value using Release semantics
void StoreRelease(hsa_signal_value_t value);
/// @brief This operation is illegal
hsa_signal_value_t WaitRelaxed(hsa_signal_condition_t condition,
hsa_signal_value_t compare_value,
uint64_t timeout, hsa_wait_state_t wait_hint) {
assert(false);
return 0;
}
/// @brief This operation is illegal
hsa_signal_value_t WaitAcquire(hsa_signal_condition_t condition,
hsa_signal_value_t compare_value,
uint64_t timeout, hsa_wait_state_t wait_hint) {
assert(false);
return 0;
}
/// @brief This operation is illegal
void AndRelaxed(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void AndAcquire(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void AndRelease(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void AndAcqRel(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void OrRelaxed(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void OrAcquire(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void OrRelease(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void OrAcqRel(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void XorRelaxed(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void XorAcquire(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void XorRelease(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void XorAcqRel(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void AddRelaxed(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void AddAcquire(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void AddRelease(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void AddAcqRel(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void SubRelaxed(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void SubAcquire(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void SubRelease(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
void SubAcqRel(hsa_signal_value_t value) { assert(false); }
/// @brief This operation is illegal
hsa_signal_value_t ExchRelaxed(hsa_signal_value_t value) {
assert(false);
return 0;
}
/// @brief This operation is illegal
hsa_signal_value_t ExchAcquire(hsa_signal_value_t value) {
assert(false);
return 0;
}
/// @brief This operation is illegal
hsa_signal_value_t ExchRelease(hsa_signal_value_t value) {
assert(false);
return 0;
}
/// @brief This operation is illegal
hsa_signal_value_t ExchAcqRel(hsa_signal_value_t value) {
assert(false);
return 0;
}
/// @brief This operation is illegal
hsa_signal_value_t CasRelaxed(hsa_signal_value_t expected,
hsa_signal_value_t value) {
assert(false);
return 0;
}
/// @brief This operation is illegal
hsa_signal_value_t CasAcquire(hsa_signal_value_t expected,
hsa_signal_value_t value) {
assert(false);
return 0;
}
/// @brief This operation is illegal
hsa_signal_value_t CasRelease(hsa_signal_value_t expected,
hsa_signal_value_t value) {
assert(false);
return 0;
}
/// @brief This operation is illegal
hsa_signal_value_t CasAcqRel(hsa_signal_value_t expected,
hsa_signal_value_t value) {
assert(false);
return 0;
}
/// @brief This operation is illegal
hsa_signal_value_t* ValueLocation() const {
assert(false);
return NULL;
}
/// @brief This operation is illegal
HsaEvent* EopEvent() {
assert(false);
return NULL;
}
// 64 byte-aligned allocation and release, for Queue::amd_queue_.
void* operator new(size_t size);
void* operator new(size_t size, void* ptr) { return ptr; }
void operator delete(void* ptr);
void operator delete(void*, void*) {}
protected:
bool _IsA(rtti_t id) const { return id == &rtti_id_; }
private:
uint32_t ComputeRingBufferMinPkts();
uint32_t ComputeRingBufferMaxPkts();
// (De)allocates and (de)registers ring_buf_.
void AllocRegisteredRingBuffer(uint32_t queue_size_pkts);
void FreeRegisteredRingBuffer();
static bool DynamicScratchHandler(hsa_signal_value_t error_code, void* arg);
// AQL packet ring buffer
void* ring_buf_;
// Size of ring_buf_ allocation.
// This may be larger than (amd_queue_.hsa_queue.size * sizeof(AqlPacket)).
uint32_t ring_buf_alloc_bytes_;
// Id of the Queue used in communication with thunk
HSA_QUEUEID queue_id_;
// Indicates is queue is valid
bool valid_;
// Indicates if queue is inactive
int32_t active_;
// Cached value of HsaNodeProperties.HSA_CAPABILITY.DoorbellType
int doorbell_type_;
// Handle of agent, which queue is attached to
GpuAgent* agent_;
hsa_profile_t agent_profile_;
uint32_t queue_full_workaround_;
// Handle of scratch memory descriptor
ScratchInfo queue_scratch_;
core::HsaEventCallback errors_callback_;
void* errors_data_;
// Is KV device queue
bool is_kv_queue_;
// Shared event used for queue errors
static HsaEvent* queue_event_;
// Queue count - used to ref count queue_event_
static volatile uint32_t queue_count_;
// Mutex for queue_event_ manipulation
static KernelMutex queue_lock_;
static int rtti_id_;
// Forbid copying and moving of this object
DISALLOW_COPY_AND_ASSIGN(AqlQueue);
};
} // namespace amd
#endif // header guard
@@ -0,0 +1,174 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef HSA_RUNTIME_CORE_INC_AMD_BLIT_KERNEL_H_
#define HSA_RUNTIME_CORE_INC_AMD_BLIT_KERNEL_H_
#include <stdint.h>
#include "core/inc/blit.h"
namespace amd {
class BlitKernel : public core::Blit {
public:
explicit BlitKernel();
virtual ~BlitKernel() override;
/// @brief Initialize a blit kernel object.
///
/// @param agent Pointer to the agent that will execute the AQL packets.
///
/// @return hsa_status_t
virtual hsa_status_t Initialize(const core::Agent& agent) override;
/// @brief Marks the blit kernel object as invalid and uncouples its link with
/// the underlying AQL kernel queue. Use of the blit object
/// once it has been release is illegal and any behavior is indeterminate
///
/// @note: The call will block until all AQL packets have been executed.
///
/// @return hsa_status_t
virtual hsa_status_t Destroy() override;
/// @brief Submit an AQL packet to perform vector copy. The call is blocking
/// until the command execution is finished.
///
/// @param dst Memory address of the copy destination.
/// @param src Memory address of the copy source.
/// @param size Size of the data to be copied.
virtual hsa_status_t SubmitLinearCopyCommand(void* dst, const void* src,
size_t size) override;
/// @brief Submit a linear copy command to the the underlying compute device's
/// control block. The call is non blocking. The memory transfer will start
/// after all dependent signals are satisfied. After the transfer is
/// completed, the out signal will be decremented.
///
/// @param dst Memory address of the copy destination.
/// @param src Memory address of the copy source.
/// @param size Size of the data to be copied.
/// @param dep_signals Arrays of dependent signal.
/// @param out_signal Output signal.
virtual hsa_status_t SubmitLinearCopyCommand(
void* dst, const void* src, size_t size,
std::vector<core::Signal*>& dep_signals,
core::Signal& out_signal) override;
/// @brief Submit an AQL packet to perform memory fill. The call is blocking
/// until the command execution is finished.
///
/// @param ptr Memory address of the fill destination.
/// @param value Value to be set.
/// @param count Number of uint32_t element to be set to the value.
virtual hsa_status_t SubmitLinearFillCommand(void* ptr, uint32_t value,
size_t count) override;
private:
struct __ALIGNED__(16) KernelCopyArgs {
const void* src;
void* dst;
uint64_t size;
uint32_t use_vector;
};
struct __ALIGNED__(16) KernelFillArgs {
void* ptr;
uint64_t num;
uint32_t value;
};
/// Reserve a slot in the queue buffer. The call will wait until the queue
/// buffer has a room.
uint64_t AcquireWriteIndex(uint32_t num_packet);
/// Update the queue doorbell register with ::write_index. This
/// function also serializes concurrent doorbell update to ensure that the
/// packet processor doesn't get invalid packet.
void ReleaseWriteIndex(uint64_t write_index, uint32_t num_packet);
/// Wait until all packets are finished.
hsa_status_t FenceRelease(uint64_t write_index, uint32_t num_copy_packet,
hsa_fence_scope_t fence);
void PopulateQueue(uint64_t index, uint64_t code_handle, void* args,
uint32_t grid_size_x, hsa_signal_t completion_signal);
KernelCopyArgs* ObtainAsyncKernelCopyArg();
/// Handles to the vector copy kernel.
uint64_t copy_code_handle_;
/// Handles to the vector copy aligned kernel.
uint64_t copy_aligned_code_handle_;
/// Handles to the fill memory kernel.
uint64_t fill_code_handle_;
/// AQL queue for submitting the vector copy kernel.
hsa_queue_t* queue_;
uint32_t queue_bitmask_;
/// Index to track concurrent kernel launch.
volatile uint64_t cached_index_;
/// Pointer to the kernel argument buffer.
void* kernarg_;
KernelCopyArgs* kernarg_async_;
uint32_t kernarg_async_mask_;
volatile uint32_t kernarg_async_counter_;
/// Completion signal for every kernel dispatched.
hsa_signal_t completion_signal_;
/// Lock to synchronize access to kernarg_ and completion_signal_
std::mutex lock_;
/// Pointer to memory containing the ISA and argument buffer.
void* code_arg_buffer_;
static const size_t kMaxCopyCount;
static const size_t kMaxFillCount;
static const uint32_t kGroupSize;
};
} // namespace amd
#endif // header guard
@@ -0,0 +1,479 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef HSA_RUNTIME_CORE_INC_AMD_BLIT_KERNEL_KV_H_
#define HSA_RUNTIME_CORE_INC_AMD_BLIT_KERNEL_KV_H_
#include <stddef.h>
#define HSA_VECTOR_COPY_KV_AKC_SIZE 368
#define HSA_VECTOR_COPY_KV_AKC_OFFSET 256
/*****HSAIL code of the ISA in ::kVectorCopyRawKv.
module &m:1:0:$full:$large:$default;
prog kernel &__vector_copy_kernel(
kernarg_u64 %src,
kernarg_u64 %dst,
kernarg_u64 %size)
{
@__vector_copy_kernel_entry:
// BB#0: // %entry
workitemabsid_u32 $s0, 0;
cvt_u64_u32 $d0, $s0;
ld_kernarg_align(8)_width(all)_u64 $d1, [%size];
cmp_ge_b1_u64 $c0, $d0, $d1;
cbr_b1 $c0, @BB0_2;
// BB#1: // %if.end
ld_kernarg_align(8)_width(all)_u64 $d1, [%src];
ld_kernarg_align(8)_width(all)_u64 $d2, [%dst];
add_u64 $d2, $d2, $d0;
add_u64 $d0, $d1, $d0;
ld_global_u8 $s0, [$d0];
st_global_u8 $s0, [$d2];
@BB0_2:
// %return
ret;
};
*/
static char kVectorCopyRawKv[] = {
127, 69, 76, 70, 2, 1, 1, 64, 0, 0, 0, 0, 0,
0, 0, 0, 1, 0, -32, 0, 1, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 64, 0, 0, 0, 0, 0, 0,
0, -104, 3, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
64, 0, 56, 0, 1, 0, 64, 0, 6, 0, 5, 0, 3,
0, 0, 96, 6, 0, 0, 0, 0, 1, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 112, 1, 0, 0, 0, 0, 0, 0,
112, 1, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0,
0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0,
1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 65, 0, -116, 0, -112, 0, 0, 0,
11, 0, 10, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 24, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 11, 0, 5, 0, 5, 0, 0, 0, 9, 0, 0,
0, 0, 0, 0, 0, 3, 0, 0, 6, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 1, 5, 0, -64, 127, 0, -116, -65,
0, -1, -128, -109, 0, 0, 16, 0, 0, 8, 0, -109, 0,
0, 0, 74, 4, 7, 64, -64, -128, 2, 2, 126, 127, 0,
-116, -65, 0, 0, -56, 125, 106, 36, -128, -66, 15, 0, -120,
-65, 0, 7, -126, -64, 127, 0, -116, -65, 4, 0, 2, 74,
5, 2, 4, 126, 2, 106, 80, -46, 2, 1, -87, 1, 0,
0, 32, -36, 1, 0, 0, 1, 6, 0, 6, 74, 7, 2,
4, 126, 4, 106, 80, -46, 2, 1, -87, 1, 112, 0, -116,
-65, 0, 0, 96, -36, 3, 1, 0, 0, 0, 0, -127, -65,
3, 0, 0, 0, 8, 0, 0, 0, 1, 0, 0, 0, 65,
77, 68, 0, 1, 0, 0, 0, 0, 0, 0, 0, 3, 0,
0, 0, 12, 0, 0, 0, 2, 0, 0, 0, 65, 77, 68,
0, 1, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 0,
3, 0, 0, 0, 28, 0, 0, 0, 3, 0, 0, 0, 65,
77, 68, 0, 4, 0, 7, 0, 7, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 65, 77, 68, 0, 65, 77, 68,
71, 80, 85, 0, 0, 3, 0, 0, 0, 40, 0, 0, 0,
4, 0, 0, 0, 65, 77, 68, 0, 26, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 65, 77, 68, 32, 72, 83,
65, 32, 82, 117, 110, 116, 105, 109, 101, 32, 70, 105, 110,
97, 108, 105, 122, 101, 114, 0, 0, 0, 38, 95, 95, 118,
101, 99, 116, 111, 114, 95, 99, 111, 112, 121, 95, 107, 101,
114, 110, 101, 108, 0, 95, 95, 104, 115, 97, 95, 115, 101,
99, 116, 105, 111, 110, 46, 104, 115, 97, 116, 101, 120, 116,
0, 0, 0, 0, 0, 0, 0, 0, 0, 26, 0, 1, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 22, 0, 0, 0, 3, 0, 1, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 46, 104, 115, 97, 116, 101, 120, 116, 0, 46, 110,
111, 116, 101, 0, 46, 115, 116, 114, 116, 97, 98, 0, 46,
115, 121, 109, 116, 97, 98, 0, 46, 115, 104, 115, 116, 114,
116, 97, 98, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0,
1, 0, 0, 0, 7, 0, -64, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0,
0, 0, 112, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 10, 0, 0, 0, 7,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 112, 2, 0, 0, 0, 0, 0,
0, -104, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 4, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 16, 0, 0, 0, 3, 0,
0, 0, 32, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 8, 3, 0, 0, 0, 0, 0, 0,
44, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 24, 0, 0, 0, 2, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 56, 3, 0, 0, 0, 0, 0, 0, 48,
0, 0, 0, 0, 0, 0, 0, 3, 0, 0, 0, 0, 0,
0, 0, 8, 0, 0, 0, 0, 0, 0, 0, 24, 0, 0,
0, 0, 0, 0, 0, 32, 0, 0, 0, 3, 0, 0, 0,
32, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 104, 3, 0, 0, 0, 0, 0, 0, 42, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0,
};
extern char* const kVectorCopyKvObject = &kVectorCopyRawKv[0];
extern size_t const kVectorCopyKvObjectSize = sizeof(kVectorCopyRawKv);
#define HSA_VECTOR_COPY_ALIGNED_KV_AKC_SIZE 436
#define HSA_VECTOR_COPY_ALIGNED_KV_AKC_OFFSET 256
/*****HSAIL code of the ISA in ::kVectorCopyAlignedRawKv.
module &m:1:0:$full:$large:$default;
extension "amd:gcn";
prog kernel &__copy_buffer_aligned_kernel(
kernarg_u64 %src,
kernarg_u64 %dst,
kernarg_u64 %size,
kernarg_u32 %use_vector)
{
@__copy_buffer_aligned_kernel_entry:
// BB#0: // %entry
workitemabsid_u32 $s0, 0;
cvt_u64_u32 $d0, $s0;
ld_kernarg_align(8)_width(all)_u64 $d1, [%size];
cmp_ge_b1_u64 $c0, $d0, $d1;
cbr_b1 $c0, @LBB0_4;
// BB#1: // %if.end
ld_kernarg_align(8)_width(all)_u64 $d2, [%dst];
ld_kernarg_align(8)_width(all)_u64 $d1, [%src];
ld_kernarg_align(4)_width(all)_u32 $s0, [%use_vector];
cmp_ne_b1_s32 $c0, $s0, 1;
cbr_b1 $c0, @LBB0_3;
// BB#2: // %if.then2
shl_u64 $d0, $d0, 4;
add_u64 $d2, $d2, $d0;
add_u64 $d0, $d1, $d0;
ld_v4_global_align(16)_const_u32 ($s0, $s1, $s2, $s3), [$d0];
st_v4_global_align(16)_u32 ($s0, $s1, $s2, $s3), [$d2];
br @LBB0_4;
@LBB0_3:
// %if.else
shl_u64 $d0, $d0, 2;
add_u64 $d2, $d2, $d0;
add_u64 $d0, $d1, $d0;
ld_global_align(4)_const_u32 $s0, [$d0];
st_global_align(4)_u32 $s0, [$d2];
@LBB0_4:
// %if.end6
ret;
};
*/
static char kVectorCopyAlignedRawKv[] = {
127, 69, 76, 70, 2, 1, 1, 64, 0, 0, 0, 0, 0,
0, 0, 0, 1, 0, -32, 0, 1, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 64, 0, 0, 0, 0, 0, 0,
0, -8, 3, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
64, 0, 56, 0, 1, 0, 64, 0, 6, 0, 5, 0, 3,
0, 0, 96, 6, 0, 0, 0, 0, 1, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, -76, 1, 0, 0, 0, 0, 0, 0,
-76, 1, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0,
1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0,
1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 65, 0, -84, 0, -112, 0, 0, 0,
11, 0, 10, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 32, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 11, 0, 7, 0, 7, 0, 0, 0, 9, 0, 0,
0, 0, 0, 0, 0, 4, 4, 4, 6, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 1, 5, 0, -64, 127, 0, -116, -65,
0, -1, -128, -109, 0, 0, 16, 0, 0, 8, 0, -109, 0,
0, 0, 74, 4, 7, 64, -64, -128, 2, 2, 126, 127, 0,
-116, -65, 0, 0, -56, 125, 106, 36, -128, -66, 32, 0, -120,
-65, 6, 7, 1, -64, 0, 7, -126, -64, 127, 0, -116, -65,
2, -127, 0, -65, 14, 0, -124, -65, 0, 0, -62, -46, 0,
9, 1, 0, 4, 0, 4, 74, 5, 2, 6, 126, 3, 3,
6, 80, 0, 0, 56, -36, 2, 0, 0, 2, 6, 0, 0,
74, 7, 2, 12, 126, 6, 3, 2, 80, 112, 0, -116, -65,
0, 0, 120, -36, 0, 2, 0, 0, 13, 0, -126, -65, 0,
0, -62, -46, 0, 5, 1, 0, 4, 0, 4, 74, 5, 2,
6, 126, 3, 3, 6, 80, 0, 0, 48, -36, 2, 0, 0,
2, 6, 0, 0, 74, 7, 2, 6, 126, 3, 3, 2, 80,
112, 0, -116, -65, 0, 0, 112, -36, 0, 2, 0, 0, 0,
0, -127, -65, 0, 0, 0, 0, 4, 0, 0, 0, 8, 0,
0, 0, 1, 0, 0, 0, 65, 77, 68, 0, 1, 0, 0,
0, 0, 0, 0, 0, 4, 0, 0, 0, 12, 0, 0, 0,
2, 0, 0, 0, 65, 77, 68, 0, 1, 0, 0, 0, 0,
0, 0, 0, 1, 1, 1, 0, 4, 0, 0, 0, 25, 0,
0, 0, 5, 0, 0, 0, 65, 77, 68, 0, 22, 0, 45,
104, 115, 97, 95, 99, 97, 108, 108, 95, 99, 111, 110, 118,
101, 110, 116, 105, 111, 110, 61, 0, 0, 0, 0, 0, 4,
0, 0, 0, 30, 0, 0, 0, 3, 0, 0, 0, 65, 77,
68, 0, 4, 0, 7, 0, 7, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 65, 77, 68, 0, 65, 77, 68, 71,
80, 85, 0, 0, 0, 0, 0, 0, 4, 0, 0, 0, 8,
0, 0, 0, 4, 0, 0, 0, 65, 77, 68, 0, -32, 101,
-118, -12, -1, 127, 0, 0, 38, 95, 95, 99, 111, 112, 121,
95, 98, 117, 102, 102, 101, 114, 95, 97, 108, 105, 103, 110,
101, 100, 95, 107, 101, 114, 110, 101, 108, 0, 95, 95, 104,
115, 97, 95, 115, 101, 99, 116, 105, 111, 110, 46, 104, 115,
97, 116, 101, 120, 116, 0, 0, 0, 0, 0, 0, 0, 0,
0, 26, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0,
-76, 1, 0, 0, 0, 0, 0, 0, 30, 0, 0, 0, 3,
0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 46, 104, 115, 97, 116, 101,
120, 116, 0, 46, 110, 111, 116, 101, 0, 46, 115, 116, 114,
116, 97, 98, 0, 46, 115, 121, 109, 116, 97, 98, 0, 46,
115, 104, 115, 116, 114, 116, 97, 98, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 1, 0, 0, 0, 1, 0, 0, 0, 7, 0, -64, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
1, 0, 0, 0, 0, 0, 0, -76, 1, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
10, 0, 0, 0, 7, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, -72, 2,
0, 0, 0, 0, 0, 0, -88, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 8, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 16,
0, 0, 0, 3, 0, 0, 0, 32, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 96, 3, 0,
0, 0, 0, 0, 0, 52, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 24, 0,
0, 0, 2, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, -104, 3, 0, 0,
0, 0, 0, 0, 48, 0, 0, 0, 0, 0, 0, 0, 3,
0, 0, 0, 0, 0, 0, 0, 8, 0, 0, 0, 0, 0,
0, 0, 24, 0, 0, 0, 0, 0, 0, 0, 32, 0, 0,
0, 3, 0, 0, 0, 32, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, -56, 3, 0, 0, 0,
0, 0, 0, 42, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0,
};
extern char* const kVectorCopyAlignedKvObject = &kVectorCopyAlignedRawKv[0];
extern size_t const kVectorCopyAlignedKvObjectSize =
sizeof(kVectorCopyAlignedRawKv);
#define HSA_FILL_MEMORY_KV_AKC_SIZE 352
#define HSA_FILL_MEMORY_KV_AKC_OFFSET 256
/*****HSAIL code of the ISA in ::kFillMemoryRawKv.
module &m:1:0:$full:$large:$default;
extension "amd:gcn";
prog kernel &__fill_memory_kernel(
kernarg_u64 %ptr,
kernarg_u64 %num,
kernarg_u32 %value)
{
@__fill_memory_kernel_entry:
// BB#0: // %entry
workitemabsid_u32 $s0, 0;
cvt_u64_u32 $d0, $s0;
ld_kernarg_align(8)_width(all)_u64 $d1, [%num];
cmp_ge_b1_u64 $c0, $d0, $d1;
cbr_b1 $c0, @LBB0_2;
// BB#1: // %if.end
ld_kernarg_align(8)_width(all)_u64 $d1, [%ptr];
ld_kernarg_align(4)_width(all)_u32 $s0, [%value];
shl_u64 $d0, $d0, 2;
add_u64 $d0, $d1, $d0;
st_global_align(4)_u32 $s0, [$d0];
@LBB0_2:
// %return
ret;
};
*/
static char kFillMemoryRawKv[] = {
127, 69, 76, 70, 2, 1, 1, 64, 0, 0, 0, 0, 0, 0,
0, 0, 1, 0, -32, 0, 1, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 64, 0, 0, 0, 0, 0, 0, 0, -104, 3,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 64, 0, 56, 0,
1, 0, 64, 0, 6, 0, 5, 0, 3, 0, 0, 96, 6, 0,
0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 96, 1,
0, 0, 0, 0, 0, 0, 96, 1, 0, 0, 0, 0, 0, 0,
0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1, 0,
0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 64, 0, -84, 0,
-112, 0, 0, 0, 11, 0, 10, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 32, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 11, 0, 3, 0, 3, 0, 0, 0, 9, 0,
0, 0, 0, 0, 0, 0, 4, 4, 4, 6, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 1, 5, 0, -64, 127, 0,
-116, -65, 0, -1, -128, -109, 0, 0, 16, 0, 0, 8, 0, -109,
0, 0, 0, 74, 2, 7, 64, -64, -128, 2, 2, 126, 127, 0,
-116, -65, 0, 0, -56, 125, 106, 36, -128, -66, 11, 0, -120, -65,
0, 7, 65, -64, 4, 7, 2, -64, 0, 0, -62, -46, 0, 5,
1, 0, 127, 0, -116, -65, 2, 0, 0, 74, 3, 2, 4, 126,
2, 3, 2, 80, 4, 2, 4, 126, 0, 0, 112, -36, 0, 2,
0, 0, 0, 0, -127, -65, 4, 0, 0, 0, 8, 0, 0, 0,
1, 0, 0, 0, 65, 77, 68, 0, 1, 0, 0, 0, 0, 0,
0, 0, 4, 0, 0, 0, 12, 0, 0, 0, 2, 0, 0, 0,
65, 77, 68, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1, 1,
1, 0, 4, 0, 0, 0, 25, 0, 0, 0, 5, 0, 0, 0,
65, 77, 68, 0, 22, 0, 45, 104, 115, 97, 95, 99, 97, 108,
108, 95, 99, 111, 110, 118, 101, 110, 116, 105, 111, 110, 61, 0,
0, 0, 0, 0, 4, 0, 0, 0, 30, 0, 0, 0, 3, 0,
0, 0, 65, 77, 68, 0, 4, 0, 7, 0, 7, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 65, 77, 68, 0, 65, 77,
68, 71, 80, 85, 0, 0, 0, 0, 0, 0, 4, 0, 0, 0,
8, 0, 0, 0, 4, 0, 0, 0, 65, 77, 68, 0, 48, 123,
44, -103, -4, 127, 0, 0, 38, 95, 95, 102, 105, 108, 108, 95,
109, 101, 109, 111, 114, 121, 95, 107, 101, 114, 110, 101, 108, 0,
95, 95, 104, 115, 97, 95, 115, 101, 99, 116, 105, 111, 110, 46,
104, 115, 97, 116, 101, 120, 116, 0, 0, 0, 0, 0, 0, 0,
0, 0, 26, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0,
96, 1, 0, 0, 0, 0, 0, 0, 22, 0, 0, 0, 3, 0,
1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 46, 104, 115, 97, 116, 101, 120, 116, 0,
46, 110, 111, 116, 101, 0, 46, 115, 116, 114, 116, 97, 98, 0,
46, 115, 121, 109, 116, 97, 98, 0, 46, 115, 104, 115, 116, 114,
116, 97, 98, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 7, 0,
-64, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 1, 0, 0, 0, 0, 0, 0, 96, 1, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 10, 0,
0, 0, 7, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 96, 2, 0, 0, 0, 0,
0, 0, -88, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 8, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 16, 0, 0, 0, 3, 0, 0, 0,
32, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 8, 3, 0, 0, 0, 0, 0, 0, 44, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
24, 0, 0, 0, 2, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 56, 3, 0, 0,
0, 0, 0, 0, 48, 0, 0, 0, 0, 0, 0, 0, 3, 0,
0, 0, 0, 0, 0, 0, 8, 0, 0, 0, 0, 0, 0, 0,
24, 0, 0, 0, 0, 0, 0, 0, 32, 0, 0, 0, 3, 0,
0, 0, 32, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 104, 3, 0, 0, 0, 0, 0, 0, 42, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0,
};
extern char* const kFillMemoryKvObject = &kFillMemoryRawKv[0];
extern size_t const kFillMemoryKvObjectSize = sizeof(kFillMemoryRawKv);
#endif // header guard
@@ -0,0 +1,490 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef HSA_RUNTIME_CORE_INC_AMD_BLIT_KERNEL_VI_H_
#define HSA_RUNTIME_CORE_INC_AMD_BLIT_KERNEL_VI_H_
#include <stddef.h>
#define HSA_VECTOR_COPY_VI_AKC_SIZE 380
#define HSA_VECTOR_COPY_VI_AKC_OFFSET 256
/*****HSAIL code of the ISA in ::kVectorCopyRawVi.
module &m:1:0:$full:$large:$default;
prog kernel &__vector_copy_kernel(
kernarg_u64 %src,
kernarg_u64 %dst,
kernarg_u64 %size)
{
@__vector_copy_kernel_entry:
// BB#0: // %entry
workitemabsid_u32 $s0, 0;
cvt_u64_u32 $d0, $s0;
ld_kernarg_align(8)_width(all)_u64 $d1, [%size];
cmp_ge_b1_u64 $c0, $d0, $d1;
cbr_b1 $c0, @BB0_2;
// BB#1: // %if.end
ld_kernarg_align(8)_width(all)_u64 $d1, [%src];
ld_kernarg_align(8)_width(all)_u64 $d2, [%dst];
add_u64 $d2, $d2, $d0;
add_u64 $d0, $d1, $d0;
ld_global_u8 $s0, [$d0];
st_global_u8 $s0, [$d2];
@BB0_2:
// %return
ret;
};
*/
static char kVectorCopyRawVi[] = {
127, 69, 76, 70, 2, 1, 1, 64, 0, 0, 0, 0, 0,
0, 0, 0, 1, 0, -32, 0, 1, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 64, 0, 0, 0, 0, 0, 0,
0, -72, 3, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
64, 0, 56, 0, 1, 0, 64, 0, 6, 0, 5, 0, 3,
0, 0, 96, 6, 0, 0, 0, 0, 1, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 124, 1, 0, 0, 0, 0, 0, 0,
124, 1, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0,
1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0,
1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, -63, 2, -84, 0, -112, 0, 0, 0,
11, 0, 10, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 32, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 96, 0, 5, 0, 5, 0, 0, 0, 9, 0, 0,
0, 0, 0, 0, 0, 4, 4, 4, 6, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 2, 0, 2, -64, 4, 0, 0, 0,
127, 0, -116, -65, 0, -1, -128, -110, 0, 0, 16, 0, 0,
8, 0, -110, 0, 0, 0, 50, 3, 0, 6, -64, 16, 0,
0, 0, -128, 2, 2, 126, 127, 0, -116, -65, 0, 0, -40,
125, 106, 32, -128, -66, 16, 0, -120, -65, 3, 1, 10, -64,
0, 0, 0, 0, 127, 0, -116, -65, 4, 0, 2, 50, 5,
2, 4, 126, 2, 106, 28, -47, 2, 1, -87, 1, 0, 0,
64, -36, 1, 0, 0, 1, 6, 0, 6, 50, 7, 2, 4,
126, 4, 106, 28, -47, 2, 1, -87, 1, 112, 0, -116, -65,
0, 0, 96, -36, 3, 1, 0, 0, 0, 0, -127, -65, 0,
0, 0, 0, 4, 0, 0, 0, 8, 0, 0, 0, 1, 0,
0, 0, 65, 77, 68, 0, 1, 0, 0, 0, 0, 0, 0,
0, 4, 0, 0, 0, 12, 0, 0, 0, 2, 0, 0, 0,
65, 77, 68, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1,
1, 1, 0, 4, 0, 0, 0, 25, 0, 0, 0, 5, 0,
0, 0, 65, 77, 68, 0, 22, 0, 45, 104, 115, 97, 95,
99, 97, 108, 108, 95, 99, 111, 110, 118, 101, 110, 116, 105,
111, 110, 61, 0, 0, 0, 0, 0, 4, 0, 0, 0, 30,
0, 0, 0, 3, 0, 0, 0, 65, 77, 68, 0, 4, 0,
7, 0, 8, 0, 0, 0, 0, 0, 0, 0, 2, 0, 0,
0, 65, 77, 68, 0, 65, 77, 68, 71, 80, 85, 0, 0,
0, 0, 0, 0, 4, 0, 0, 0, 8, 0, 0, 0, 4,
0, 0, 0, 65, 77, 68, 0, 32, 103, -72, 81, -3, 127,
0, 0, 38, 95, 95, 118, 101, 99, 116, 111, 114, 95, 99,
111, 112, 121, 95, 107, 101, 114, 110, 101, 108, 0, 95, 95,
104, 115, 97, 95, 115, 101, 99, 116, 105, 111, 110, 46, 104,
115, 97, 116, 101, 120, 116, 0, 0, 0, 0, 0, 0, 0,
0, 0, 26, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0,
0, 124, 1, 0, 0, 0, 0, 0, 0, 22, 0, 0, 0,
3, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 46, 104, 115, 97, 116,
101, 120, 116, 0, 46, 110, 111, 116, 101, 0, 46, 115, 116,
114, 116, 97, 98, 0, 46, 115, 121, 109, 116, 97, 98, 0,
46, 115, 104, 115, 116, 114, 116, 97, 98, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 7, 0, -64,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 1, 0, 0, 0, 0, 0, 0, 124, 1, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 10, 0, 0, 0, 7, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, -128,
2, 0, 0, 0, 0, 0, 0, -88, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 8, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
16, 0, 0, 0, 3, 0, 0, 0, 32, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 40, 3,
0, 0, 0, 0, 0, 0, 44, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 24,
0, 0, 0, 2, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 88, 3, 0,
0, 0, 0, 0, 0, 48, 0, 0, 0, 0, 0, 0, 0,
3, 0, 0, 0, 0, 0, 0, 0, 8, 0, 0, 0, 0,
0, 0, 0, 24, 0, 0, 0, 0, 0, 0, 0, 32, 0,
0, 0, 3, 0, 0, 0, 32, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, -120, 3, 0, 0,
0, 0, 0, 0, 42, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
};
extern char* const kVectorCopyViObject = &kVectorCopyRawVi[0];
extern size_t const kVectorCopyViObjectSize = sizeof(kVectorCopyRawVi);
#define HSA_VECTOR_COPY_ALIGNED_VI_AKC_SIZE 452
#define HSA_VECTOR_COPY_ALIGNED_VI_AKC_OFFSET 256
/*****HSAIL code of the ISA in ::kVectorCopyAlignedRawVi.
module &m:1:0:$full:$large:$default;
extension "amd:gcn";
prog kernel &__copy_buffer_aligned_kernel(
kernarg_u64 %src,
kernarg_u64 %dst,
kernarg_u64 %size,
kernarg_u32 %use_vector)
{
@__copy_buffer_aligned_kernel_entry:
// BB#0: // %entry
workitemabsid_u32 $s0, 0;
cvt_u64_u32 $d0, $s0;
ld_kernarg_align(8)_width(all)_u64 $d1, [%size];
cmp_ge_b1_u64 $c0, $d0, $d1;
cbr_b1 $c0, @LBB0_4;
// BB#1: // %if.end
ld_kernarg_align(8)_width(all)_u64 $d2, [%dst];
ld_kernarg_align(8)_width(all)_u64 $d1, [%src];
ld_kernarg_align(4)_width(all)_u32 $s0, [%use_vector];
cmp_ne_b1_s32 $c0, $s0, 1;
cbr_b1 $c0, @LBB0_3;
// BB#2: // %if.then2
shl_u64 $d0, $d0, 4;
add_u64 $d2, $d2, $d0;
add_u64 $d0, $d1, $d0;
ld_v4_global_align(16)_const_u32 ($s0, $s1, $s2, $s3), [$d0];
st_v4_global_align(16)_u32 ($s0, $s1, $s2, $s3), [$d2];
br @LBB0_4;
@LBB0_3:
// %if.else
shl_u64 $d0, $d0, 2;
add_u64 $d2, $d2, $d0;
add_u64 $d0, $d1, $d0;
ld_global_align(4)_const_u32 $s0, [$d0];
st_global_align(4)_u32 $s0, [$d2];
@LBB0_4:
// %if.end6
ret;
};
*/
static char kVectorCopyAlignedRawVi[] = {
127, 69, 76, 70, 2, 1, 1, 64, 0, 0, 0, 0, 0,
0, 0, 0, 1, 0, -32, 0, 1, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 64, 0, 0, 0, 0, 0, 0,
0, 8, 4, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
64, 0, 56, 0, 1, 0, 64, 0, 6, 0, 5, 0, 3,
0, 0, 96, 6, 0, 0, 0, 0, 1, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, -60, 1, 0, 0, 0, 0, 0, 0,
-60, 1, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0,
1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0,
1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 65, 0, -84, 0, -112, 0, 0, 0,
11, 0, 74, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 32, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 16, 0, 8, 0, 8, 0, 0, 0, 12, 0, 0,
0, 0, 0, 0, 0, 4, 4, 4, 6, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 2, 0, 2, -64, 4, 0, 0, 0,
127, 0, -116, -65, 0, -1, -128, -110, 0, 0, 16, 0, 0,
8, 0, -110, 0, 0, 0, 50, 3, 0, 6, -64, 16, 0,
0, 0, -128, 2, 2, 126, 127, 0, -116, -65, 0, 0, -40,
125, 106, 32, -128, -66, 34, 0, -120, -65, -125, 0, 2, -64,
24, 0, 0, 0, 3, 2, 10, -64, 0, 0, 0, 0, 127,
0, -116, -65, 2, -127, 0, -65, 14, 0, -124, -65, 0, 0,
-113, -46, -124, 0, 2, 0, 8, 0, 4, 50, 9, 2, 6,
126, 3, 3, 6, 56, 0, 0, 92, -36, 2, 0, 0, 4,
10, 0, 0, 50, 11, 2, 4, 126, 2, 3, 2, 56, 112,
0, -116, -65, 0, 0, 124, -36, 0, 4, 0, 0, 13, 0,
-126, -65, 0, 0, -113, -46, -126, 0, 2, 0, 8, 0, 4,
50, 9, 2, 6, 126, 3, 3, 6, 56, 0, 0, 80, -36,
2, 0, 0, 4, 10, 0, 0, 50, 11, 2, 4, 126, 2,
3, 2, 56, 112, 0, -116, -65, 0, 0, 112, -36, 0, 4,
0, 0, 0, 0, -127, -65, 0, 0, 0, 0, 4, 0, 0,
0, 8, 0, 0, 0, 1, 0, 0, 0, 65, 77, 68, 0,
1, 0, 0, 0, 0, 0, 0, 0, 4, 0, 0, 0, 12,
0, 0, 0, 2, 0, 0, 0, 65, 77, 68, 0, 1, 0,
0, 0, 0, 0, 0, 0, 1, 1, 1, 0, 4, 0, 0,
0, 25, 0, 0, 0, 5, 0, 0, 0, 65, 77, 68, 0,
22, 0, 45, 104, 115, 97, 95, 99, 97, 108, 108, 95, 99,
111, 110, 118, 101, 110, 116, 105, 111, 110, 61, 0, 0, 0,
0, 0, 4, 0, 0, 0, 30, 0, 0, 0, 3, 0, 0,
0, 65, 77, 68, 0, 4, 0, 7, 0, 8, 0, 0, 0,
0, 0, 0, 0, 1, 0, 0, 0, 65, 77, 68, 0, 65,
77, 68, 71, 80, 85, 0, 0, 0, 0, 0, 0, 4, 0,
0, 0, 8, 0, 0, 0, 4, 0, 0, 0, 65, 77, 68,
0, 96, 62, -27, 85, -1, 127, 0, 0, 38, 95, 95, 99,
111, 112, 121, 95, 98, 117, 102, 102, 101, 114, 95, 97, 108,
105, 103, 110, 101, 100, 95, 107, 101, 114, 110, 101, 108, 0,
95, 95, 104, 115, 97, 95, 115, 101, 99, 116, 105, 111, 110,
46, 104, 115, 97, 116, 101, 120, 116, 0, 0, 0, 0, 0,
0, 0, 0, 0, 26, 0, 1, 0, 0, 0, 0, 0, 0,
0, 0, 0, -60, 1, 0, 0, 0, 0, 0, 0, 30, 0,
0, 0, 3, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 46, 104, 115,
97, 116, 101, 120, 116, 0, 46, 110, 111, 116, 101, 0, 46,
115, 116, 114, 116, 97, 98, 0, 46, 115, 121, 109, 116, 97,
98, 0, 46, 115, 104, 115, 116, 114, 116, 97, 98, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 7,
0, -64, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 1, 0, 0, 0, 0, 0, 0, -60, 1, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 10, 0, 0, 0, 7, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, -56, 2, 0, 0, 0, 0, 0, 0, -88, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 8,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 16, 0, 0, 0, 3, 0, 0, 0, 32, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
112, 3, 0, 0, 0, 0, 0, 0, 52, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 24, 0, 0, 0, 2, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, -88,
3, 0, 0, 0, 0, 0, 0, 48, 0, 0, 0, 0, 0,
0, 0, 3, 0, 0, 0, 0, 0, 0, 0, 8, 0, 0,
0, 0, 0, 0, 0, 24, 0, 0, 0, 0, 0, 0, 0,
32, 0, 0, 0, 3, 0, 0, 0, 32, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, -40, 3,
0, 0, 0, 0, 0, 0, 42, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
};
extern char* const kVectorCopyAlignedViObject = &kVectorCopyAlignedRawVi[0];
extern size_t const kVectorCopyAlignedViObjectSize =
sizeof(kVectorCopyAlignedRawVi);
#define HSA_FILL_MEMORY_VI_AKC_SIZE 368
#define HSA_FILL_MEMORY_VI_AKC_OFFSET 256
/*****HSAIL code of the ISA in ::kFillMemoryRawVi.
module &m:1:0:$full:$large:$default;
extension "amd:gcn";
prog kernel &__fill_memory_kernel(
kernarg_u64 %ptr,
kernarg_u64 %num,
kernarg_u32 %value)
{
@__fill_memory_kernel_entry:
// BB#0: // %entry
workitemabsid_u32 $s0, 0;
cvt_u64_u32 $d0, $s0;
ld_kernarg_align(8)_width(all)_u64 $d1, [%num];
cmp_ge_b1_u64 $c0, $d0, $d1;
cbr_b1 $c0, @LBB0_2;
// BB#1: // %if.end
ld_kernarg_align(8)_width(all)_u64 $d1, [%ptr];
ld_kernarg_align(4)_width(all)_u32 $s0, [%value];
shl_u64 $d0, $d0, 2;
add_u64 $d0, $d1, $d0;
st_global_align(4)_u32 $s0, [$d0];
@LBB0_2:
// %return
ret;
};
*/
static char kFillMemoryRawVi[] = {
127, 69, 76, 70, 2, 1, 1, 64, 0, 0, 0, 0, 0,
0, 0, 0, 1, 0, -32, 0, 1, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 64, 0, 0, 0, 0, 0, 0,
0, -88, 3, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
64, 0, 56, 0, 1, 0, 64, 0, 6, 0, 5, 0, 3,
0, 0, 96, 6, 0, 0, 0, 0, 1, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 112, 1, 0, 0, 0, 0, 0, 0,
112, 1, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0,
1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0,
1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 64, 0, -84, 0, -112, 0, 0, 0,
11, 0, 74, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 32, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 13, 0, 3, 0, 3, 0, 0, 0, 9, 0, 0,
0, 0, 0, 0, 0, 4, 4, 4, 6, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 2, 0, 2, -64, 4, 0, 0, 0,
127, 0, -116, -65, 0, -1, -128, -110, 0, 0, 16, 0, 0,
8, 0, -110, 0, 0, 0, 50, 3, 0, 6, -64, 8, 0,
0, 0, -128, 2, 2, 126, 127, 0, -116, -65, 0, 0, -40,
125, 106, 32, -128, -66, 13, 0, -120, -65, -125, 0, 6, -64,
0, 0, 0, 0, 3, 1, 2, -64, 16, 0, 0, 0, 0,
0, -113, -46, -126, 0, 2, 0, 127, 0, -116, -65, 2, 0,
0, 50, 3, 2, 4, 126, 2, 3, 2, 56, 4, 2, 4,
126, 0, 0, 112, -36, 0, 2, 0, 0, 0, 0, -127, -65,
4, 0, 0, 0, 8, 0, 0, 0, 1, 0, 0, 0, 65,
77, 68, 0, 1, 0, 0, 0, 0, 0, 0, 0, 4, 0,
0, 0, 12, 0, 0, 0, 2, 0, 0, 0, 65, 77, 68,
0, 1, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 0,
4, 0, 0, 0, 25, 0, 0, 0, 5, 0, 0, 0, 65,
77, 68, 0, 22, 0, 45, 104, 115, 97, 95, 99, 97, 108,
108, 95, 99, 111, 110, 118, 101, 110, 116, 105, 111, 110, 61,
0, 0, 0, 0, 0, 4, 0, 0, 0, 30, 0, 0, 0,
3, 0, 0, 0, 65, 77, 68, 0, 4, 0, 7, 0, 8,
0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 65, 77,
68, 0, 65, 77, 68, 71, 80, 85, 0, 0, 0, 0, 0,
0, 4, 0, 0, 0, 8, 0, 0, 0, 4, 0, 0, 0,
65, 77, 68, 0, 16, -20, 88, 97, -4, 127, 0, 0, 38,
95, 95, 102, 105, 108, 108, 95, 109, 101, 109, 111, 114, 121,
95, 107, 101, 114, 110, 101, 108, 0, 95, 95, 104, 115, 97,
95, 115, 101, 99, 116, 105, 111, 110, 46, 104, 115, 97, 116,
101, 120, 116, 0, 0, 0, 0, 0, 0, 0, 0, 0, 26,
0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 112, 1,
0, 0, 0, 0, 0, 0, 22, 0, 0, 0, 3, 0, 1,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 46, 104, 115, 97, 116, 101, 120, 116,
0, 46, 110, 111, 116, 101, 0, 46, 115, 116, 114, 116, 97,
98, 0, 46, 115, 121, 109, 116, 97, 98, 0, 46, 115, 104,
115, 116, 114, 116, 97, 98, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1,
0, 0, 0, 1, 0, 0, 0, 7, 0, -64, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0,
0, 0, 0, 0, 0, 112, 1, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 10, 0,
0, 0, 7, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 112, 2, 0, 0,
0, 0, 0, 0, -88, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 8, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 16, 0, 0,
0, 3, 0, 0, 0, 32, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 24, 3, 0, 0, 0,
0, 0, 0, 44, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 0, 0, 24, 0, 0, 0,
2, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0, 72, 3, 0, 0, 0, 0,
0, 0, 48, 0, 0, 0, 0, 0, 0, 0, 3, 0, 0,
0, 0, 0, 0, 0, 8, 0, 0, 0, 0, 0, 0, 0,
24, 0, 0, 0, 0, 0, 0, 0, 32, 0, 0, 0, 3,
0, 0, 0, 32, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 120, 3, 0, 0, 0, 0, 0,
0, 42, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0,
0, 0, 0, 0, 0, 0, 0,
};
extern char* const kFillMemoryViObject = &kFillMemoryRawVi[0];
extern size_t const kFillMemoryViObjectSize = sizeof(kFillMemoryRawVi);
#endif // header guard
@@ -0,0 +1,218 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef HSA_RUNTIME_CORE_INC_AMD_BLIT_SDMA_H_
#define HSA_RUNTIME_CORE_INC_AMD_BLIT_SDMA_H_
#include <stdint.h>
#include "hsakmt.h"
#include "core/inc/blit.h"
#include "core/inc/runtime.h"
#include "core/inc/signal.h"
#include "core/util/utils.h"
namespace amd {
class BlitSdma : public core::Blit {
public:
explicit BlitSdma();
virtual ~BlitSdma() override;
/// @brief Initialize a User Mode SDMA Queue object. Input parameters specify
/// properties of queue being created.
///
/// @param agent Pointer to the agent that will execute the PM4 commands.
///
/// @return hsa_status_t
virtual hsa_status_t Initialize(const core::Agent& agent) override;
/// @brief Marks the queue object as invalid and uncouples its link with
/// the underlying compute device's control block. Use of queue object
/// once it has been release is illegal and any behavior is indeterminate
///
/// @note: The call will block until all packets have executed.
///
/// @return hsa_status_t
virtual hsa_status_t Destroy() override;
/// @brief Submit a linear copy command to the queue buffer.
///
/// @param dst Memory address of the copy destination.
/// @param src Memory address of the copy source.
/// @param size Size of the data to be copied.
virtual hsa_status_t SubmitLinearCopyCommand(void* dst, const void* src,
size_t size) override;
/// @brief Submit a linear copy command to the the underlying compute device's
/// control block. The call is non blocking. The memory transfer will start
/// after all dependent signals are satisfied. After the transfer is
/// completed, the out signal will be decremented.
///
/// @param dst Memory address of the copy destination.
/// @param src Memory address of the copy source.
/// @param size Size of the data to be copied.
/// @param dep_signals Arrays of dependent signal.
/// @param out_signal Output signal.
virtual hsa_status_t SubmitLinearCopyCommand(
void* dst, const void* src, size_t size,
std::vector<core::Signal*>& dep_signals,
core::Signal& out_signal) override;
/// @brief Submit a linear fill command to the queue buffer
///
/// @param ptr Memory address of the fill destination.
/// @param value Value to be set.
/// @param count Number of uint32_t element to be set to the value.
virtual hsa_status_t SubmitLinearFillCommand(void* ptr, uint32_t value,
size_t count) override;
protected:
/// @brief Acquires the address into queue buffer where a new command
/// packet of specified size could be written. The address that is
/// returned is guaranteed to be unique even in a multi-threaded access
/// scenario. This function is guaranteed to return a pointer for writing
/// data into the queue buffer.
///
/// @param cmd_size Command packet size in bytes.
///
/// @return pointer into the queue buffer where a PM4 packet of specified size
/// could be written. NULL if input size is greater than the size of queue
/// buffer.
char* AcquireWriteAddress(uint32_t cmd_size);
void UpdateWriteAndDoorbellRegister(uint32_t current_offset,
uint32_t new_offset);
/// @brief Updates the Write Register of compute device to the end of
/// SDMA packet written into queue buffer. The update to Write Register
/// will be safe under multi-threaded usage scenario. Furthermore, updates
/// to Write Register are blocking until all prior updates are completed
/// i.e. if two threads T1 & T2 were to call release, then updates by T2
/// will block until T1 has completed its update (assumes T1 acquired the
/// write address first).
///
/// @param cmd_addr pointer into the queue buffer where a PM4 packet was
/// written.
///
/// @param cmd_size Command packet size in bytes.
void ReleaseWriteAddress(char* cmd_addr, uint32_t cmd_size);
/// @brief Writes NO-OP words into queue buffer in case writing a command
/// causes the queue buffer to wrap.
///
/// @param cmd_size Size in bytes of command causing queue buffer to wrap.
void WrapQueue(uint32_t cmd_size);
/// @brief Build fence command
void BuildFenceCommand(char* fence_command_addr, uint32_t* fence,
uint32_t fence_value);
uint32_t* ObtainFenceObject();
void WaitFence(uint32_t* fence, uint32_t fence_value);
void BuildCopyCommand(char* cmd_addr, uint32_t num_copy_command, void* dst,
const void* src, size_t size);
void BuildPollCommand(char* cmd_addr, void* addr, uint32_t reference);
void BuildAtomicDecrementCommand(char* cmd_addr, void* addr);
/// Indicates size of Queue buffer in bytes.
uint32_t queue_size_;
/// Base address of the Queue buffer at construction time.
char* queue_start_addr_;
uint32_t* fence_base_addr_;
uint32_t fence_pool_size_;
uint32_t fence_pool_mask_;
volatile uint32_t fence_pool_counter_;
/// Queue resource descriptor for doorbell, read
/// and write indices
HsaQueueResource queue_resource_;
/// @brief Current address of execution in Queue buffer.
///
/// @note: The value of address is obtained by reading
/// the value of Write Register of the compute device.
/// Users should write to the Queue buffer at the current
/// address, else it will lead to execution error and potentially
/// a hang.
///
/// @note: The value of Write Register does not always begin
/// with Zero after a Queue has been created. This needs to be
/// understood better. This means that current address number of
/// words of Queue buffer is unavailable for use.
volatile uint32_t cached_reserve_offset_;
volatile uint32_t cached_commit_offset_;
uint32_t linear_copy_command_size_;
uint32_t fill_command_size_;
uint32_t fence_command_size_;
uint32_t poll_command_size_;
uint32_t atomic_command_size_;
// Max copy size of a single linear copy command packet.
size_t max_single_linear_copy_size_;
/// Max total copy size supported by the queue.
size_t max_total_linear_copy_size_;
/// Max count of uint32_t of a single fill command packet.
size_t max_single_fill_size_;
/// Max total fill count supported by the queue.
size_t max_total_fill_size_;
std::mutex wrap_lock_;
};
} // namespace amd
#endif // header guard
@@ -0,0 +1,154 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
// AMD specific HSA backend.
#ifndef HSA_RUNTIME_CORE_INC_AMD_CPU_AGENT_H_
#define HSA_RUNTIME_CORE_INC_AMD_CPU_AGENT_H_
#include <vector>
#include "hsakmt.h"
#include "core/inc/runtime.h"
#include "core/inc/agent.h"
#include "core/inc/queue.h"
namespace amd {
// @brief Class to represent a CPU device.
class CpuAgent : public core::Agent {
public:
// @brief CpuAgent constructor.
//
// @param [in] node Node id. Each CPU in different socket will get distinct
// id.
// @param [in] node_props Node property.
CpuAgent(HSAuint32 node, const HsaNodeProperties& node_props);
// @brief CpuAgent destructor.
~CpuAgent();
// @brief Invoke the user provided callback for each region accessible by
// this agent.
//
// @param [in] include_peer If true, the callback will be also invoked on each
// peer memory region accessible by this agent. If false, only invoke the
// callback on memory region owned by this agent.
// @param [in] callback User provided callback function.
// @param [in] data User provided pointer as input for @p callback.
//
// @retval ::HSA_STATUS_SUCCESS if the callback function for each traversed
// region returns ::HSA_STATUS_SUCCESS.
hsa_status_t VisitRegion(bool include_peer,
hsa_status_t (*callback)(hsa_region_t region,
void* data),
void* data) const;
// @brief Override from core::Agent.
hsa_status_t IterateRegion(hsa_status_t (*callback)(hsa_region_t region,
void* data),
void* data) const override;
// @brief Override from core::Agent.
hsa_status_t GetInfo(hsa_agent_info_t attribute, void* value) const override;
// @brief Override from core::Agent.
hsa_status_t QueueCreate(size_t size, hsa_queue_type_t queue_type,
core::HsaEventCallback event_callback, void* data,
uint32_t private_segment_size,
uint32_t group_segment_size,
core::Queue** queue) override;
// @brief Returns number of data caches.
__forceinline size_t num_cache() const { return cache_props_.size(); }
// @brief Returns data cache property.
//
// @param [in] idx Cache level.
__forceinline const HsaCacheProperties& cache_prop(int idx) const {
return cache_props_[idx];
}
// @brief Override from core::Agent.
const std::vector<const core::MemoryRegion*>& regions() const override {
return regions_;
}
// @brief OVerride from core::Agent.
const core::Isa* isa() const override { return NULL; }
private:
// @brief Query the driver to get the region list owned by this agent.
void InitRegionList();
// @brief Query the driver to get the cache properties.
void InitCacheList();
// @brief Invoke the user provided callback for every region in @p regions.
//
// @param [in] regions Array of region object.
// @param [in] callback User provided callback function.
// @param [in] data User provided pointer as input for @p callback.
//
// @retval ::HSA_STATUS_SUCCESS if the callback function for each traversed
// region returns ::HSA_STATUS_SUCCESS.
hsa_status_t VisitRegion(
const std::vector<const core::MemoryRegion*>& regions,
hsa_status_t (*callback)(hsa_region_t region, void* data),
void* data) const;
// @brief Node property.
const HsaNodeProperties properties_;
// @brief Array of data cache property. The array index represents the cache
// level.
std::vector<HsaCacheProperties> cache_props_;
// @brief Array of regions owned by this agent.
std::vector<const core::MemoryRegion*> regions_;
DISALLOW_COPY_AND_ASSIGN(CpuAgent);
};
} // namespace amd
#endif // header guard
@@ -0,0 +1,222 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef AMD_ELF_IMAGE_HPP_
#define AMD_ELF_IMAGE_HPP_
#include <string>
#include <sstream>
#include <vector>
#include <memory>
namespace amd {
namespace elf {
class Symbol;
class SymbolTable;
class Section;
class RelocationSection;
class Segment {
public:
virtual ~Segment() { }
virtual uint64_t type() const = 0;
virtual uint64_t memSize() const = 0;
virtual uint64_t align() const = 0;
virtual uint64_t imageSize() const = 0;
virtual uint64_t vaddr() const = 0;
virtual uint64_t flags() const = 0;
virtual const char* data() const = 0;
virtual uint16_t getSegmentIndex() = 0;
virtual bool updateAddSection(Section *section) = 0;
};
class Section {
public:
virtual ~Section() { }
virtual uint16_t getSectionIndex() const = 0;
virtual uint32_t type() const = 0;
virtual std::string Name() const = 0;
virtual uint64_t offset() const = 0;
virtual uint64_t addr() const = 0;
virtual bool updateAddr(uint64_t addr) = 0;
virtual uint64_t addralign() const = 0;
virtual uint64_t flags() const = 0;
virtual uint64_t size() const = 0;
virtual uint64_t nextDataOffset(uint64_t align) const = 0;
virtual uint64_t addData(const void *src, uint64_t size, uint64_t align) = 0;
virtual bool getData(uint64_t offset, void* dest, uint64_t size) = 0;
virtual Segment* segment() = 0;
virtual RelocationSection* asRelocationSection() = 0;
virtual bool hasRelocationSection() const = 0;
virtual RelocationSection* relocationSection(SymbolTable* symtab = 0) = 0;
virtual bool setMemSize(uint64_t s) = 0;
virtual uint64_t memSize() const = 0;
virtual bool setAlign(uint64_t a) = 0;
virtual uint64_t memAlign() const = 0;
};
class Relocation {
public:
virtual ~Relocation() { }
virtual RelocationSection* section() = 0;
virtual uint32_t type() = 0;
virtual uint32_t symbolIndex() = 0;
virtual Symbol* symbol() = 0;
virtual uint64_t offset() = 0;
virtual int64_t addend() = 0;
};
class RelocationSection : public virtual Section {
public:
virtual Relocation* addRelocation(uint32_t type, Symbol* symbol, uint64_t offset, int64_t addend) = 0;
virtual size_t relocationCount() const = 0;
virtual Relocation* relocation(size_t i) = 0;
virtual Section* targetSection() = 0;
};
class StringTable : public virtual Section {
public:
virtual const char* addString(const std::string& s) = 0;
virtual size_t addString1(const std::string& s) = 0;
virtual const char* getString(size_t ndx) = 0;
virtual size_t getStringIndex(const char* name) = 0;
};
class Symbol {
public:
virtual ~Symbol() { }
virtual uint32_t index() = 0;
virtual uint32_t type() = 0;
virtual uint32_t binding() = 0;
virtual uint64_t size() = 0;
virtual uint64_t value() = 0;
virtual unsigned char other() = 0;
virtual std::string name() = 0;
virtual Section* section() = 0;
virtual void setValue(uint64_t value) = 0;
virtual void setSize(uint64_t size) = 0;
};
class SymbolTable : public virtual Section {
public:
virtual Symbol* addSymbol(Section* section, const std::string& name, uint64_t value, uint64_t size, unsigned char type, unsigned char binding, unsigned char other = 0) = 0;
virtual size_t symbolCount() = 0;
virtual Symbol* symbol(size_t i) = 0;
};
class NoteSection : public virtual Section {
public:
virtual bool addNote(const std::string& name, uint32_t type, const void* desc = 0, uint32_t desc_size = 0) = 0;
virtual bool getNote(const std::string& name, uint32_t type, void** desc, uint32_t* desc_size) = 0;
};
class Image {
public:
virtual ~Image() { }
virtual bool initNew(uint16_t machine, uint16_t type, uint8_t os_abi = 0, uint8_t abi_version = 0, uint32_t e_flags = 0) = 0;
virtual bool loadFromFile(const std::string& filename) = 0;
virtual bool saveToFile(const std::string& filename) = 0;
virtual bool initFromBuffer(const void* buffer, size_t size) = 0;
virtual bool initAsBuffer(const void* buffer, size_t size) = 0;
virtual bool writeTo(const std::string& filename) = 0;
virtual bool copyToBuffer(void** buf, size_t* size = 0) = 0; // Copy to new buffer allocated with malloc
virtual bool copyToBuffer(void* buf, size_t size) = 0; // Copy to existing buffer of given size.
virtual const char* data() = 0;
virtual uint64_t size() = 0;
virtual uint16_t Machine() = 0;
virtual uint16_t Type() = 0;
std::string output() { return out.str(); }
virtual bool Freeze() = 0;
virtual bool Validate() = 0;
virtual StringTable* shstrtab() = 0;
virtual StringTable* strtab() = 0;
virtual SymbolTable* symtab() = 0;
virtual SymbolTable* getSymtab(uint16_t index) = 0;
virtual StringTable* addStringTable(const std::string& name) = 0;
virtual StringTable* getStringTable(uint16_t index) = 0;
virtual SymbolTable* addSymbolTable(const std::string& name, StringTable* stab = 0) = 0;
virtual size_t segmentCount() = 0;
virtual Segment* segment(size_t i) = 0;
virtual Segment* segmentByVAddr(uint64_t vaddr) = 0;
virtual size_t sectionCount() = 0;
virtual Section* section(size_t i) = 0;
virtual Section* sectionByVAddr(uint64_t vaddr) = 0;
virtual NoteSection* note() = 0;
virtual NoteSection* addNoteSection(const std::string& name) = 0;
virtual Segment* initSegment(uint32_t type, uint32_t flags, uint64_t paddr = 0) = 0;
virtual bool addSegments() = 0;
virtual Section* addSection(const std::string &name,
uint32_t type,
uint64_t flags = 0,
uint64_t entsize = 0,
Segment* segment = 0) = 0;
virtual RelocationSection* relocationSection(Section* sec, SymbolTable* symtab = 0) = 0;
protected:
std::ostringstream out;
};
Image* NewElf32Image();
Image* NewElf64Image();
uint64_t ElfSize(const void* buffer);
std::string GetNoteString(uint32_t s_size, const char* s);
}
}
#endif // AMD_ELF_IMAGE_HPP_
@@ -0,0 +1,354 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
// AMD specific HSA backend.
#ifndef HSA_RUNTIME_CORE_INC_AMD_GPU_AGENT_H_
#define HSA_RUNTIME_CORE_INC_AMD_GPU_AGENT_H_
#include <vector>
#include "hsakmt.h"
#include "core/inc/runtime.h"
#include "core/inc/agent.h"
#include "core/inc/blit.h"
#include "core/inc/signal.h"
#include "core/util/small_heap.h"
#include "core/util/locks.h"
namespace amd {
// @brief Contains scratch memory information.
struct ScratchInfo {
void* queue_base;
size_t size;
size_t size_per_thread;
ptrdiff_t queue_process_offset;
};
// @brief Interface to represent a GPU agent.
class GpuAgentInt : public core::Agent {
public:
// @brief Constructor
GpuAgentInt(uint32_t node_id)
: core::Agent(node_id, core::Agent::DeviceType::kAmdGpuDevice) {}
// @brief Invoke the user provided callback for each region accessible by
// this agent.
//
// @param [in] include_peer If true, the callback will be also invoked on each
// peer memory region accessible by this agent. If false, only invoke the
// callback on memory region owned by this agent.
// @param [in] callback User provided callback function.
// @param [in] data User provided pointer as input for @p callback.
//
// @retval ::HSA_STATUS_SUCCESS if the callback function for each traversed
// region returns ::HSA_STATUS_SUCCESS.
virtual hsa_status_t VisitRegion(bool include_peer,
hsa_status_t (*callback)(hsa_region_t region,
void* data),
void* data) const = 0;
// @brief Carve scratch memory from scratch pool.
//
// @param [out] scratch Structure to be populated with the carved memory
// information.
virtual void AcquireQueueScratch(ScratchInfo& scratch) = 0;
// @brief Release scratch memory back to scratch pool.
//
// @param [in] base Address of scratch memory previously acquired with
// call to ::AcquireQueueScratch.
virtual void ReleaseQueueScratch(void* base) = 0;
// @brief Translate the kernel start and end dispatch timestamp from agent
// domain to host domain.
//
// @param [in] signal Pointer to signal that provides the dispatch timing.
// @param [out] time Structure to be populated with the host domain value.
virtual void TranslateTime(core::Signal* signal,
hsa_amd_profiling_dispatch_time_t& time) = 0;
// @brief Translate timestamp agent domain to host domain.
//
// @param [out] time Timestamp in agent domain.
virtual uint64_t TranslateTime(uint64_t tick) = 0;
// @brief Sets the coherency type of this agent.
//
// @param [in] type New coherency type.
//
// @retval true The new coherency type is set successfuly.
virtual bool current_coherency_type(hsa_amd_coherency_type_t type) = 0;
// @brief Returns the current coherency type of this agent.
//
// @retval Coherency type.
virtual hsa_amd_coherency_type_t current_coherency_type() const = 0;
// @brief Query if agent represent Kaveri GPU.
//
// @retval true if agent is Kaveri GPU.
virtual bool is_kv_device() const = 0;
// @brief Query the agent HSA profile.
//
// @retval HSA profile.
virtual hsa_profile_t profile() const = 0;
};
class GpuAgent : public GpuAgentInt {
public:
// @brief GPU agent constructor.
//
// @param [in] node Node id. Each CPU in different socket will get distinct
// id.
// @param [in] node_props Node property.
GpuAgent(HSAuint32 node, const HsaNodeProperties& node_props);
// @brief GPU agent destructor.
~GpuAgent();
// @brief Initialize DMA queue.
//
// @retval HSA_STATUS_SUCCESS DMA queue initialization is successful.
hsa_status_t InitDma();
uint16_t GetMicrocodeVersion() const;
// @brief Assembles SP3 shader source into executable code.
//
// @param [in] src_sp3 SP3 shader source text representation.
// @param [in] func_name Name of the SP3 function to assemble.
// @param [out] code_buf Executable code buffer.
// @param [out] code_buf_size Size of executable code buffer in bytes.
void AssembleShader(const char* src_sp3, const char* func_name,
void*& code_buf, size_t& code_buf_size);
// @brief Frees executable code created by AssembleShader.
//
// @param [in] code_buf Executable code buffer.
// @param [in] code_buf_size Size of executable code buffer in bytes.
void ReleaseShader(void* code_buf, size_t code_buf_size);
// @brief Override from core::Agent.
hsa_status_t VisitRegion(bool include_peer,
hsa_status_t (*callback)(hsa_region_t region,
void* data),
void* data) const override;
// @brief Override from core::Agent.
hsa_status_t IterateRegion(hsa_status_t (*callback)(hsa_region_t region,
void* data),
void* data) const override;
// @brief Override from core::Agent.
hsa_status_t DmaCopy(void* dst, const void* src, size_t size) override;
// @brief Override from core::Agent.
hsa_status_t DmaCopy(void* dst, const void* src, size_t size,
std::vector<core::Signal*>& dep_signals,
core::Signal& out_signal) override;
// @brief Override from core::Agent.
hsa_status_t DmaFill(void* ptr, uint32_t value, size_t count) override;
// @brief Override from core::Agent.
hsa_status_t GetInfo(hsa_agent_info_t attribute, void* value) const override;
// @brief Override from core::Agent.
hsa_status_t QueueCreate(size_t size, hsa_queue_type_t queue_type,
core::HsaEventCallback event_callback, void* data,
uint32_t private_segment_size,
uint32_t group_segment_size,
core::Queue** queue) override;
// @brief Override from amd::GpuAgentInt.
void AcquireQueueScratch(ScratchInfo& scratch) override;
// @brief Override from amd::GpuAgentInt.
void ReleaseQueueScratch(void* base) override;
// @brief Override from amd::GpuAgentInt.
void TranslateTime(core::Signal* signal,
hsa_amd_profiling_dispatch_time_t& time) override;
// @brief Override from amd::GpuAgentInt.
uint64_t TranslateTime(uint64_t tick) override;
// @brief Override from amd::GpuAgentInt.
bool current_coherency_type(hsa_amd_coherency_type_t type) override;
// @brief Override from amd::GpuAgentInt.
hsa_amd_coherency_type_t current_coherency_type() const override {
return current_coherency_type_;
}
// Getter & setters.
// @brief Returns node property.
__forceinline const HsaNodeProperties& properties() const {
return properties_;
}
// @brief Returns number of data caches.
__forceinline size_t num_cache() const { return cache_props_.size(); }
// @brief Returns data cache property.
//
// @param [in] idx Cache level.
__forceinline const HsaCacheProperties& cache_prop(int idx) const {
return cache_props_[idx];
}
// @brief Override from core::Agent.
const std::vector<const core::MemoryRegion*>& regions() const override {
return regions_;
}
// @brief OVerride from core::Agent.
const core::Isa* isa() const override { return isa_; }
// @brief Override from amd::GpuAgentInt.
__forceinline bool is_kv_device() const override { return is_kv_device_; }
// @brief Override from amd::GpuAgentInt.
__forceinline hsa_profile_t profile() const override { return profile_; }
protected:
static const uint32_t minAqlSize_ = 0x1000; // 4KB min
static const uint32_t maxAqlSize_ = 0x20000; // 8MB max
// @brief Invoke the user provided callback for every region in @p regions.
//
// @param [in] regions Array of region object.
// @param [in] callback User provided callback function.
// @param [in] data User provided pointer as input for @p callback.
//
// @retval ::HSA_STATUS_SUCCESS if the callback function for each traversed
// region returns ::HSA_STATUS_SUCCESS.
hsa_status_t VisitRegion(
const std::vector<const core::MemoryRegion*>& regions,
hsa_status_t (*callback)(hsa_region_t region, void* data),
void* data) const;
// @brief Update ::t1_ tick count.
void SyncClocks();
// @brief Binds the second-level trap handler to this node.
void BindTrapHandler();
// @brief Node properties.
const HsaNodeProperties properties_;
// @brief Current coherency type.
hsa_amd_coherency_type_t current_coherency_type_;
// @brief Maximum number of queues that can be created.
uint32_t max_queues_;
// @brief Object to manage scratch memory.
SmallHeap scratch_pool_;
// @brief Default scratch size per queue.
size_t queue_scratch_len_;
// @brief Default scratch size per work item.
size_t scratch_per_thread_;
// @brief Blit object to handle memory copy/fill.
core::Blit* blit_;
// @brief Mutex to protect the update to coherency type.
KernelMutex coherency_lock_;
// @brief Mutex to protect access to scratch pool.
KernelMutex scratch_lock_;
// @brief Mutex to protect access to ::t1_.
KernelMutex t1_lock_;
// @brief GPU tick on initialization.
HsaClockCounters t0_;
HsaClockCounters t1_;
// @brief Array of GPU cache property.
std::vector<HsaCacheProperties> cache_props_;
// @brief Array of regions owned by this agent.
std::vector<const core::MemoryRegion*> regions_;
core::Isa* isa_;
// @brief HSA profile.
hsa_profile_t profile_;
bool is_kv_device_;
void* trap_code_buf_;
size_t trap_code_buf_size_;
private:
// @brief Query the driver to get the region list owned by this agent.
void InitRegionList();
// @brief Reserve memory for scratch pool to be used by AQL queue of this
// agent.
void InitScratchPool();
// @brief Query the driver to get the cache properties.
void InitCacheList();
// @brief Alternative aperture base address. Only on KV.
uintptr_t ape1_base_;
// @brief Alternative aperture size. Only on KV.
size_t ape1_size_;
DISALLOW_COPY_AND_ASSIGN(GpuAgent);
};
} // namespace
#endif // header guard
@@ -0,0 +1,387 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef AMD_HSA_CODE_HPP_
#define AMD_HSA_CODE_HPP_
#include "amd_elf_image.hpp"
#include "amd_hsa_elf.h"
#include "amd_hsa_kernel_code.h"
#include "hsa.h"
#include "hsa_ext_finalize.h"
#include <memory>
#include <sstream>
#include <cassert>
#include <unordered_map>
namespace amd {
namespace hsa {
namespace common {
template<uint64_t signature>
class Signed {
public:
static const uint64_t CT_SIGNATURE;
const uint64_t RT_SIGNATURE;
protected:
Signed(): RT_SIGNATURE(signature) {}
virtual ~Signed() {}
};
template<uint64_t signature>
const uint64_t Signed<signature>::CT_SIGNATURE = signature;
bool IsAccessibleMemoryAddress(uint64_t address);
template<typename class_type, typename member_type>
size_t OffsetOf(member_type class_type::*member)
{
return (char*)&((class_type*)nullptr->*member) - (char*)nullptr;
}
template<typename class_type>
class_type* ObjectAt(uint64_t address)
{
if (!IsAccessibleMemoryAddress(address)) {
return nullptr;
}
const uint64_t *rt_signature =
(const uint64_t*)(address + OffsetOf(&class_type::RT_SIGNATURE));
if (nullptr == rt_signature) {
return nullptr;
}
if (class_type::CT_SIGNATURE != *rt_signature) {
return nullptr;
}
return (class_type*)address;
}
}
namespace code {
typedef amd::elf::Segment Segment;
typedef amd::elf::Section Section;
typedef amd::elf::RelocationSection RelocationSection;
typedef amd::elf::Relocation Relocation;
class KernelSymbol;
class VariableSymbol;
class Symbol {
protected:
amd::elf::Symbol* elfsym;
public:
explicit Symbol(amd::elf::Symbol* elfsym_)
: elfsym(elfsym_) { }
virtual ~Symbol() { }
virtual bool IsKernelSymbol() const { return false; }
virtual KernelSymbol* AsKernelSymbol() { assert(false); return 0; }
virtual bool IsVariableSymbol() const { return false; }
virtual VariableSymbol* AsVariableSymbol() { assert(false); return 0; }
amd::elf::Symbol* elfSym() { return elfsym; }
std::string Name() const { return elfsym ? elfsym->name() : ""; }
Section* GetSection() { return elfsym->section(); }
virtual uint64_t SectionOffset() const { return elfsym->value(); }
virtual uint64_t VAddr() const { return elfsym->section()->addr() + elfsym->value(); }
uint32_t Index() const { return elfsym ? elfsym->index() : 0; }
bool IsDeclaration() const;
bool IsDefinition() const;
virtual bool IsAgent() const;
virtual hsa_symbol_kind_t Kind() const = 0;
hsa_symbol_linkage_t Linkage() const;
hsa_variable_allocation_t Allocation() const;
hsa_variable_segment_t Segment() const;
uint64_t Size() const;
uint32_t Size32() const;
uint32_t Alignment() const;
bool IsConst() const;
virtual hsa_status_t GetInfo(hsa_code_symbol_info_t attribute, void *value);
static hsa_code_symbol_t ToHandle(Symbol* sym);
static Symbol* FromHandle(hsa_code_symbol_t handle);
void setValue(uint64_t value) { elfsym->setValue(value); }
void setSize(uint32_t size) { elfsym->setSize(size); }
};
class KernelSymbol : public Symbol {
private:
uint32_t kernarg_segment_size, kernarg_segment_alignment;
uint32_t group_segment_size, private_segment_size;
bool is_dynamic_callstack;
public:
explicit KernelSymbol(amd::elf::Symbol* elfsym_, const amd_kernel_code_t* akc);
bool IsKernelSymbol() const override { return true; }
KernelSymbol* AsKernelSymbol() override { return this; }
hsa_symbol_kind_t Kind() const override { return HSA_SYMBOL_KIND_KERNEL; }
hsa_status_t GetInfo(hsa_code_symbol_info_t attribute, void *value) override;
};
class VariableSymbol : public Symbol {
public:
explicit VariableSymbol(amd::elf::Symbol* elfsym_)
: Symbol(elfsym_) { }
bool IsVariableSymbol() const override { return true; }
VariableSymbol* AsVariableSymbol() override { return this; }
hsa_symbol_kind_t Kind() const override { return HSA_SYMBOL_KIND_VARIABLE; }
hsa_status_t GetInfo(hsa_code_symbol_info_t attribute, void *value) override;
};
class AmdHsaCode {
private:
std::ostringstream out;
std::unique_ptr<amd::elf::Image> img;
std::vector<Segment*> dataSegments;
std::vector<Section*> dataSections;
std::vector<RelocationSection*> relocationSections;
std::vector<Symbol*> symbols;
bool combineDataSegments;
Segment* hsaSegments[AMDGPU_HSA_SEGMENT_LAST][2];
Section* hsaSections[AMDGPU_HSA_SECTION_LAST];
amd::elf::Section* hsatext;
amd::elf::Section* imageInit;
amd::elf::Section* samplerInit;
amd::elf::Section* debugInfo;
amd::elf::Section* debugLine;
amd::elf::Section* debugAbbrev;
bool PullElf();
bool PullElfV1();
bool PullElfV2();
void AddAmdNote(uint32_t type, const void* desc, uint32_t desc_size);
template <typename S>
bool GetAmdNote(uint32_t type, S** desc)
{
uint32_t desc_size;
if (!img->note()->getNote("AMD", type, (void**) desc, &desc_size)) {
out << "Failed to find note, type: " << type << std::endl;
return false;
}
if (desc_size < sizeof(S)) {
out << "Note size mismatch, type: " << type << " size: " << desc_size << " expected at least " << sizeof(S) << std::endl;
return false;
}
return true;
}
void PrintSegment(std::ostream& out, Segment* segment);
void PrintSection(std::ostream& out, Section* section);
void PrintRawData(std::ostream& out, Section* section);
void PrintRawData(std::ostream& out, const unsigned char *data, size_t size);
void PrintRelocationData(std::ostream& out, RelocationSection* section);
void PrintSymbol(std::ostream& out, Symbol* sym);
void PrintDisassembly(std::ostream& out, const unsigned char *isa, size_t size, uint32_t isa_offset = 0);
std::string MangleSymbolName(const std::string& module_name, const std::string symbol_name);
bool ElfImageError();
public:
bool HasHsaText() const { return hsatext != 0; }
amd::elf::Section* HsaText() { assert(hsatext); return hsatext; }
const amd::elf::Section* HsaText() const { assert(hsatext); return hsatext; }
amd::elf::SymbolTable* Symtab() { assert(img); return img->symtab(); }
uint16_t Machine() { return img->Machine(); }
AmdHsaCode(bool combineDataSegments = true);
virtual ~AmdHsaCode();
std::string output() { return out.str(); }
bool LoadFromFile(const std::string& filename);
bool SaveToFile(const std::string& filename);
bool WriteToBuffer(void* buffer);
bool InitFromBuffer(const void* buffer, size_t size);
bool InitAsBuffer(const void* buffer, size_t size);
bool InitAsHandle(hsa_code_object_t code_handle);
bool InitNew(bool xnack = false);
bool Freeze();
hsa_code_object_t GetHandle();
const char* ElfData();
uint64_t ElfSize();
bool Validate();
void Print(std::ostream& out);
void PrintNotes(std::ostream& out);
void PrintSegments(std::ostream& out);
void PrintSections(std::ostream& out);
void PrintSymbols(std::ostream& out);
void PrintMachineCode(std::ostream& out);
void PrintMachineCode(std::ostream& out, KernelSymbol* sym);
bool PrintToFile(const std::string& filename);
void AddNoteCodeObjectVersion(uint32_t major, uint32_t minor);
bool GetNoteCodeObjectVersion(uint32_t* major, uint32_t* minor);
bool GetNoteCodeObjectVersion(std::string& version);
void AddNoteHsail(uint32_t hsail_major, uint32_t hsail_minor, hsa_profile_t profile, hsa_machine_model_t machine_model, hsa_default_float_rounding_mode_t rounding_mode);
bool GetNoteHsail(uint32_t* hsail_major, uint32_t* hsail_minor, hsa_profile_t* profile, hsa_machine_model_t* machine_model, hsa_default_float_rounding_mode_t* default_float_round);
void AddNoteIsa(const std::string& vendor_name, const std::string& architecture_name, uint32_t major, uint32_t minor, uint32_t stepping);
bool GetNoteIsa(std::string& vendor_name, std::string& architecture_name, uint32_t* major_version, uint32_t* minor_version, uint32_t* stepping);
bool GetNoteIsa(std::string& isaName);
void AddNoteProducer(uint32_t major, uint32_t minor, const std::string& producer);
bool GetNoteProducer(uint32_t* major, uint32_t* minor, std::string& producer_name);
void AddNoteProducerOptions(const std::string& options);
void AddNoteProducerOptions(int32_t call_convention, const hsa_ext_control_directives_t& user_directives, const std::string& user_options);
bool GetNoteProducerOptions(std::string& options);
hsa_status_t GetInfo(hsa_code_object_info_t attribute, void *value);
hsa_status_t GetSymbol(const char *module_name, const char *symbol_name, hsa_code_symbol_t *sym);
hsa_status_t IterateSymbols(hsa_code_object_t code_object,
hsa_status_t (*callback)(
hsa_code_object_t code_object,
hsa_code_symbol_t symbol,
void* data),
void* data);
void AddHsaTextData(const void* buffer, size_t size);
uint64_t NextKernelCodeOffset() const;
bool AddKernelCode(KernelSymbol* sym, const void* code, size_t size);
Symbol* AddKernelDefinition(const std::string& name, const void* isa, size_t isa_size);
size_t DataSegmentCount() { return dataSegments.size(); }
Segment* DataSegment(size_t i) { return dataSegments[i]; }
size_t DataSectionCount() { return dataSections.size(); }
Section* DataSection(size_t i) { return dataSections[i]; }
Section* AddEmptySection();
Section* AddCodeSection(Segment* segment);
Section* AddDataSection(const std::string &name,
uint32_t type,
uint64_t flags,
Segment* segment);
bool HasImageInitSection() const { return imageInit != 0; }
Section* ImageInitSection();
void AddImageInitializer(Symbol* image, uint64_t destOffset, const amdgpu_hsa_image_descriptor_t& init);
void AddImageInitializer(Symbol* image, uint64_t destOffset,
amdgpu_hsa_metadata_kind16_t kind,
amdgpu_hsa_image_geometry8_t geometry,
amdgpu_hsa_image_channel_order8_t channel_order, amdgpu_hsa_image_channel_type8_t channel_type,
uint64_t width, uint64_t height, uint64_t depth, uint64_t array);
bool HasSamplerInitSection() const { return samplerInit != 0; }
amd::elf::Section* SamplerInitSection();
amd::elf::Section* AddSamplerInit();
void AddSamplerInitializer(Symbol* sampler, uint64_t destOffset, const amdgpu_hsa_sampler_descriptor_t& init);
void AddSamplerInitializer(Symbol* sampler, uint64_t destOffset,
amdgpu_hsa_sampler_coord8_t coord,
amdgpu_hsa_sampler_filter8_t filter,
amdgpu_hsa_sampler_addressing8_t addressing);
void AddInitVarWithAddress(bool large, Symbol* dest, uint64_t destOffset, Symbol* addrOf, uint64_t addrAddend);
void InitHsaSegment(amdgpu_hsa_elf_segment_t segment, bool writable);
bool AddHsaSegments();
Segment* HsaSegment(amdgpu_hsa_elf_segment_t segment, bool writable);
void InitHsaSectionSegment(amdgpu_hsa_elf_section_t section, bool combineSegments = true);
Section* HsaDataSection(amdgpu_hsa_elf_section_t section, bool combineSegments = true);
Symbol* AddExecutableSymbol(const std::string &name,
unsigned char type,
unsigned char binding,
unsigned char other,
Section *section = 0);
Symbol* AddVariableSymbol(const std::string &name,
unsigned char type,
unsigned char binding,
unsigned char other,
Section *section,
uint64_t value,
uint64_t size);
void AddSectionSymbols();
size_t RelocationSectionCount() { return relocationSections.size(); }
RelocationSection* GetRelocationSection(size_t i) { return relocationSections[i]; }
size_t SymbolCount() { return symbols.size(); }
Symbol* GetSymbol(size_t i) { return symbols[i]; }
Symbol* GetSymbolByElfIndex(size_t index);
Symbol* FindSymbol(const std::string &n);
void AddData(amdgpu_hsa_elf_section_t section, const void* data = 0, size_t size = 0);
Section* DebugInfo();
Section* DebugLine();
Section* DebugAbbrev();
Section* AddHsaHlDebug(const std::string& name, const void* data, size_t size);
};
class AmdHsaCodeManager {
private:
typedef std::unordered_map<uint64_t, AmdHsaCode*> CodeMap;
CodeMap codeMap;
public:
AmdHsaCode* FromHandle(hsa_code_object_t handle);
bool Destroy(hsa_code_object_t handle);
};
class KernelSymbolV2 : public KernelSymbol {
private:
public:
explicit KernelSymbolV2(amd::elf::Symbol* elfsym_, const amd_kernel_code_t* akc);
bool IsAgent() const override { return true; }
uint64_t SectionOffset() const override { return elfsym->value() - elfsym->section()->addr(); }
uint64_t VAddr() const override { return elfsym->value(); }
};
class VariableSymbolV2 : public VariableSymbol {
private:
public:
explicit VariableSymbolV2(amd::elf::Symbol* elfsym_) : VariableSymbol(elfsym_) { }
bool IsAgent() const override { return false; }
uint64_t SectionOffset() const override { return elfsym->value() - elfsym->section()->addr(); }
uint64_t VAddr() const override { return elfsym->value(); }
};
}
}
}
#endif // AMD_HSA_CODE_HPP_
@@ -0,0 +1,358 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef AMD_HSA_LOADER_HPP
#define AMD_HSA_LOADER_HPP
#include <cstddef>
#include <cstdint>
#include "hsa.h"
#include "hsa_ext_image.h"
#include "amd_hsa_elf.h"
#include "amd_load_map.h"
#include <string>
#include <mutex>
#include <vector>
/// @brief Major version of the AMD HSA Loader. Major versions are not backwards
/// compatible.
#define AMD_HSA_LOADER_VERSION_MAJOR 0
/// @brief Minor version of the AMD HSA Loader. Minor versions are backwards
/// compatible.
#define AMD_HSA_LOADER_VERSION_MINOR 5
/// @brief Descriptive version of the AMD HSA Loader.
#define AMD_HSA_LOADER_VERSION "AMD HSA Loader v0.05 (June 16, 2015)"
enum hsa_ext_symbol_info_t {
HSA_EXT_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT_SIZE = 100,
HSA_EXT_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT_ALIGN = 101,
};
typedef uint32_t hsa_symbol_info32_t;
typedef hsa_executable_symbol_t hsa_symbol_t;
typedef hsa_executable_symbol_info_t hsa_symbol_info_t;
namespace amd {
namespace hsa {
namespace loader {
//===----------------------------------------------------------------------===//
// Context. //
//===----------------------------------------------------------------------===//
class Context {
public:
virtual ~Context() {}
virtual hsa_isa_t IsaFromName(const char *name) = 0;
virtual bool IsaSupportedByAgent(hsa_agent_t agent, hsa_isa_t isa) = 0;
virtual void* SegmentAlloc(amdgpu_hsa_elf_segment_t segment, hsa_agent_t agent, size_t size, size_t align, bool zero) = 0;
virtual bool SegmentCopy(amdgpu_hsa_elf_segment_t segment, hsa_agent_t agent, void* dst, size_t offset, const void* src, size_t size) = 0;
virtual void SegmentFree(amdgpu_hsa_elf_segment_t segment, hsa_agent_t agent, void* seg, size_t size) = 0;
virtual void* SegmentAddress(amdgpu_hsa_elf_segment_t segment, hsa_agent_t agent, void* seg, size_t offset) = 0;
virtual void* SegmentHostAddress(amdgpu_hsa_elf_segment_t segment, hsa_agent_t agent, void* seg, size_t offset) = 0;
virtual bool SegmentFreeze(amdgpu_hsa_elf_segment_t segment, hsa_agent_t agent, void* seg, size_t size) = 0;
virtual bool ImageExtensionSupported() = 0;
virtual hsa_status_t ImageCreate(
hsa_agent_t agent,
hsa_access_permission_t image_permission,
const hsa_ext_image_descriptor_t *image_descriptor,
const void *image_data,
hsa_ext_image_t *image_handle) = 0;
virtual hsa_status_t ImageDestroy(
hsa_agent_t agent, hsa_ext_image_t image_handle) = 0;
virtual hsa_status_t SamplerCreate(
hsa_agent_t agent,
const hsa_ext_sampler_descriptor_t *sampler_descriptor,
hsa_ext_sampler_t *sampler_handle) = 0;
virtual hsa_status_t SamplerDestroy(
hsa_agent_t agent, hsa_ext_sampler_t sampler_handle) = 0;
protected:
Context() {}
private:
Context(const Context &c);
Context& operator=(const Context &c);
};
//===----------------------------------------------------------------------===//
// Symbol. //
//===----------------------------------------------------------------------===//
class Symbol {
public:
static hsa_symbol_t Handle(Symbol *symbol) {
hsa_symbol_t symbol_handle =
{reinterpret_cast<uint64_t>(symbol)};
return symbol_handle;
}
static Symbol* Object(hsa_symbol_t symbol_handle) {
Symbol *symbol =
reinterpret_cast<Symbol*>(symbol_handle.handle);
return symbol;
}
virtual ~Symbol() {}
virtual bool GetInfo(hsa_symbol_info32_t symbol_info, void *value) = 0;
protected:
Symbol() {}
private:
Symbol(const Symbol &s);
Symbol& operator=(const Symbol &s);
};
//===----------------------------------------------------------------------===//
// LoadedCodeObject. //
//===----------------------------------------------------------------------===//
class LoadedCodeObject {
public:
static amd_loaded_code_object_t Handle(LoadedCodeObject *object) {
amd_loaded_code_object_t handle =
{reinterpret_cast<uint64_t>(object)};
return handle;
}
static LoadedCodeObject* Object(amd_loaded_code_object_t handle) {
LoadedCodeObject *object =
reinterpret_cast<LoadedCodeObject*>(handle.handle);
return object;
}
virtual ~LoadedCodeObject() {}
virtual bool GetInfo(amd_loaded_code_object_info_t attribute, void *value) = 0;
virtual hsa_status_t IterateLoadedSegments(
hsa_status_t (*callback)(
amd_loaded_segment_t loaded_segment,
void *data),
void *data) = 0;
protected:
LoadedCodeObject() {}
private:
LoadedCodeObject(const LoadedCodeObject&);
LoadedCodeObject& operator=(const LoadedCodeObject&);
};
//===----------------------------------------------------------------------===//
// LoadedSegment. //
//===----------------------------------------------------------------------===//
class LoadedSegment {
public:
static amd_loaded_segment_t Handle(LoadedSegment *object) {
amd_loaded_segment_t handle =
{reinterpret_cast<uint64_t>(object)};
return handle;
}
static LoadedSegment* Object(amd_loaded_segment_t handle) {
LoadedSegment *object =
reinterpret_cast<LoadedSegment*>(handle.handle);
return object;
}
virtual ~LoadedSegment() {}
virtual bool GetInfo(amd_loaded_segment_info_t attribute, void *value) = 0;
protected:
LoadedSegment() {}
private:
LoadedSegment(const LoadedSegment&);
LoadedSegment& operator=(const LoadedSegment&);
};
//===----------------------------------------------------------------------===//
// Executable. //
//===----------------------------------------------------------------------===//
class Executable {
public:
static hsa_executable_t Handle(Executable *executable) {
hsa_executable_t executable_handle =
{reinterpret_cast<uint64_t>(executable)};
return executable_handle;
}
static Executable* Object(hsa_executable_t executable_handle) {
Executable *executable =
reinterpret_cast<Executable*>(executable_handle.handle);
return executable;
}
virtual ~Executable() {}
virtual hsa_status_t GetInfo(
hsa_executable_info_t executable_info, void *value) = 0;
virtual hsa_status_t DefineProgramExternalVariable(
const char *name, void *address) = 0;
virtual hsa_status_t DefineAgentExternalVariable(
const char *name,
hsa_agent_t agent,
hsa_variable_segment_t segment,
void *address) = 0;
virtual hsa_status_t LoadCodeObject(
hsa_agent_t agent,
hsa_code_object_t code_object,
const char *options,
amd_loaded_code_object_t *loaded_code_object = nullptr) = 0;
virtual hsa_status_t LoadCodeObject(
hsa_agent_t agent,
hsa_code_object_t code_object,
size_t code_object_size,
const char *options,
amd_loaded_code_object_t *loaded_code_object = nullptr) = 0;
virtual hsa_status_t Freeze(const char *options) = 0;
virtual hsa_status_t Validate(uint32_t *result) = 0;
virtual Symbol* GetSymbol(
const char *module_name,
const char *symbol_name,
hsa_agent_t agent,
int32_t call_convention) = 0;
typedef hsa_status_t (*iterate_symbols_f)(
hsa_executable_t executable,
hsa_symbol_t symbol_handle,
void *data);
virtual hsa_status_t IterateSymbols(
iterate_symbols_f callback, void *data) = 0;
virtual hsa_status_t IterateLoadedCodeObjects(
hsa_status_t (*callback)(
amd_loaded_code_object_t loaded_code_object,
void *data),
void *data) = 0;
protected:
Executable() {}
private:
Executable(const Executable &e);
Executable& operator=(const Executable &e);
static std::vector<Executable*> executables;
static std::mutex executables_mutex;
};
/// @class Loader
class Loader {
public:
/// @brief Destructor.
virtual ~Loader() {}
/// @brief Creates AMD HSA Loader with specified @p context.
///
/// @param[in] context Context. Must not be null.
///
/// @returns AMD HSA Loader on success, null on failure.
static Loader* Create(Context* context);
/// @brief Destroys AMD HSA Loader @p Loader_object.
///
/// @param[in] loader AMD HSA Loader to destroy. Must not be null.
static void Destroy(Loader *loader);
/// @returns Context associated with Loader.
virtual Context* GetContext() const = 0;
/// @brief Creates empty AMD HSA Executable with specified @p profile,
/// @p options
virtual Executable* CreateExecutable(hsa_profile_t profile, const char *options) = 0;
virtual void DestroyExecutable(Executable *executable) = 0;
virtual hsa_status_t IterateExecutables(
hsa_status_t (*callback)(
hsa_executable_t executable,
void *data),
void *data) = 0;
protected:
/// @brief Default constructor.
Loader() {}
private:
/// @brief Copy constructor - not available.
Loader(const Loader&);
/// @brief Assignment operator - not available.
Loader& operator=(const Loader&);
};
} // namespace loader
} // namespace hsa
} // namespace amd
#endif // AMD_HSA_LOADER_HPP
+174
View File
@@ -0,0 +1,174 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef AMD_LOAD_MAP_H
#define AMD_LOAD_MAP_H
#include "hsa.h"
#ifdef __cplusplus
extern "C" {
#endif // __cplusplus
/// @todo.
enum {
AMD_EXTENSION_LOAD_MAP = 0x1002
};
/// @todo.
typedef struct amd_loaded_code_object_s {
uint64_t handle;
} amd_loaded_code_object_t;
/// @todo.
enum amd_loaded_code_object_info_t {
AMD_LOADED_CODE_OBJECT_INFO_ELF_IMAGE = 0,
AMD_LOADED_CODE_OBJECT_INFO_ELF_IMAGE_SIZE = 1
};
/// @todo.
typedef struct amd_loaded_segment_s {
uint64_t handle;
} amd_loaded_segment_t;
/// @todo.
enum amd_loaded_segment_info_t {
AMD_LOADED_SEGMENT_INFO_TYPE = 0,
AMD_LOADED_SEGMENT_INFO_ELF_BASE_ADDRESS = 1,
AMD_LOADED_SEGMENT_INFO_LOAD_BASE_ADDRESS = 2,
AMD_LOADED_SEGMENT_INFO_SIZE = 3
};
/// @todo.
hsa_status_t amd_executable_load_code_object(
hsa_executable_t executable,
hsa_agent_t agent,
hsa_code_object_t code_object,
const char *options,
amd_loaded_code_object_t *loaded_code_object);
/// @brief Invokes @p callback for each available executable in current
/// process.
hsa_status_t amd_iterate_executables(
hsa_status_t (*callback)(
hsa_executable_t executable,
void *data),
void *data);
/// @brief Invokes @p callback for each loaded code object in specified
/// @p executable.
hsa_status_t amd_executable_iterate_loaded_code_objects(
hsa_executable_t executable,
hsa_status_t (*callback)(
amd_loaded_code_object_t loaded_code_object,
void *data),
void *data);
/// @brief Retrieves current value of specified @p loaded_code_object's
/// @p attribute.
hsa_status_t amd_loaded_code_object_get_info(
amd_loaded_code_object_t loaded_code_object,
amd_loaded_code_object_info_t attribute,
void *value);
/// @brief Invokes @p callback for each loaded segment in specified
/// @p loaded_code_object.
hsa_status_t amd_loaded_code_object_iterate_loaded_segments(
amd_loaded_code_object_t loaded_code_object,
hsa_status_t (*callback)(
amd_loaded_segment_t loaded_segment,
void *data),
void *data);
/// @brief Retrieves current value of specified @p loaded_segment's
/// @p attribute.
hsa_status_t amd_loaded_segment_get_info(
amd_loaded_segment_t loaded_segment,
amd_loaded_segment_info_t attribute,
void *value);
#define amd_load_map_1_00
typedef struct amd_load_map_1_00_pfn_s {
hsa_status_t (*amd_executable_load_code_object)(
hsa_executable_t executable,
hsa_agent_t agent,
hsa_code_object_t code_object,
const char *options,
amd_loaded_code_object_t *loaded_code_object);
hsa_status_t (*amd_iterate_executables)(
hsa_status_t (*callback)(
hsa_executable_t executable,
void *data),
void *data);
hsa_status_t (*amd_executable_iterate_loaded_code_objects)(
hsa_executable_t executable,
hsa_status_t (*callback)(
amd_loaded_code_object_t loaded_code_object,
void *data),
void *data);
hsa_status_t (*amd_loaded_code_object_get_info)(
amd_loaded_code_object_t loaded_code_object,
amd_loaded_code_object_info_t attribute,
void *value);
hsa_status_t (*amd_loaded_code_object_iterate_loaded_segments)(
amd_loaded_code_object_t loaded_code_object,
hsa_status_t (*callback)(
amd_loaded_segment_t loaded_segment,
void *data),
void *data);
hsa_status_t (*amd_loaded_segment_get_info)(
amd_loaded_segment_t loaded_segment,
amd_loaded_segment_info_t attribute,
void *value);
} amd_load_map_1_00_pfn_t;
#ifdef __cplusplus
} // extern "C"
#endif // __cplusplus
#endif // AMD_LOAD_MAP_H
@@ -0,0 +1,97 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef HSA_RUNTIME_CORE_INC_AMD_LOADER_CONTEXT_HPP
#define HSA_RUNTIME_CORE_INC_AMD_LOADER_CONTEXT_HPP
#include "core/inc/amd_hsa_loader.hpp"
namespace amd {
class LoaderContext final: public hsa::loader::Context {
public:
LoaderContext(): hsa::loader::Context() {}
~LoaderContext() {}
hsa_isa_t IsaFromName(const char *name) override;
bool IsaSupportedByAgent(hsa_agent_t agent, hsa_isa_t code_object_isa) override;
void* SegmentAlloc(amdgpu_hsa_elf_segment_t segment, hsa_agent_t agent, size_t size, size_t align, bool zero) override;
bool SegmentCopy(amdgpu_hsa_elf_segment_t segment, hsa_agent_t agent, void* dst, size_t offset, const void* src, size_t size) override;
void SegmentFree(amdgpu_hsa_elf_segment_t segment, hsa_agent_t agent, void* seg, size_t size = 0) override;
void* SegmentAddress(amdgpu_hsa_elf_segment_t segment, hsa_agent_t agent, void* seg, size_t offset) override;
void* SegmentHostAddress(amdgpu_hsa_elf_segment_t segment, hsa_agent_t agent, void* seg, size_t offset) override;
bool SegmentFreeze(amdgpu_hsa_elf_segment_t segment, hsa_agent_t agent, void* seg, size_t size) override;
bool ImageExtensionSupported();
hsa_status_t ImageCreate(
hsa_agent_t agent,
hsa_access_permission_t image_permission,
const hsa_ext_image_descriptor_t *image_descriptor,
const void *image_data,
hsa_ext_image_t *image_handle);
hsa_status_t ImageDestroy(hsa_agent_t agent, hsa_ext_image_t image_handle);
hsa_status_t SamplerCreate(
hsa_agent_t agent,
const hsa_ext_sampler_descriptor_t *sampler_descriptor,
hsa_ext_sampler_t *sampler_handle);
hsa_status_t SamplerDestroy(hsa_agent_t agent, hsa_ext_sampler_t sampler_handle);
private:
LoaderContext(const LoaderContext&);
LoaderContext& operator=(const LoaderContext&);
};
} // namespace amd
#endif // HSA_RUNTIME_CORE_INC_AMD_LOADER_CONTEXT_HPP
@@ -0,0 +1,191 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
// AMD specific HSA backend.
#ifndef HSA_RUNTIME_CORE_INC_AMD_MEMORY_REGION_H_
#define HSA_RUNTIME_CORE_INC_AMD_MEMORY_REGION_H_
#include "hsakmt.h"
#include "core/inc/agent.h"
#include "core/inc/memory_region.h"
#include "inc/hsa_ext_amd.h"
namespace amd {
class MemoryRegion : public core::MemoryRegion {
public:
/// @brief Convert this object into hsa_region_t.
static __forceinline hsa_region_t Convert(MemoryRegion* region) {
const hsa_region_t region_handle = {
static_cast<uint64_t>(reinterpret_cast<uintptr_t>(region))};
return region_handle;
}
static __forceinline const hsa_region_t Convert(const MemoryRegion* region) {
const hsa_region_t region_handle = {
static_cast<uint64_t>(reinterpret_cast<uintptr_t>(region))};
return region_handle;
}
/// @brief Convert hsa_region_t into amd::MemoryRegion *.
static __forceinline MemoryRegion* Convert(hsa_region_t region) {
return reinterpret_cast<MemoryRegion*>(region.handle);
}
/// @brief Allocate agent accessible memory (system / local memory).
static void* AllocateKfdMemory(const HsaMemFlags& flag, HSAuint32 node_id,
size_t size);
/// @brief Free agent accessible memory (system / local memory).
static void FreeKfdMemory(void* ptr, size_t size);
static bool RegisterMemory(void* ptr, size_t size, size_t num_nodes,
const uint32_t* nodes);
static void DeregisterMemory(void* ptr);
/// @brief Pin memory.
static bool MakeKfdMemoryResident(size_t num_node, const uint32_t* nodes,
void* ptr, size_t size,
uint64_t* alternate_va,
HsaMemMapFlags map_flag);
/// @brief Unpin memory.
static void MakeKfdMemoryUnresident(void* ptr);
MemoryRegion(bool fine_grain, bool full_profile, core::Agent* owner,
const HsaMemoryProperties& mem_props);
~MemoryRegion();
hsa_status_t Allocate(size_t size, void** address) const;
hsa_status_t Allocate(bool restrict_access, size_t size,
void** address) const;
hsa_status_t Free(void* address, size_t size) const;
hsa_status_t GetInfo(hsa_region_info_t attribute, void* value) const;
hsa_status_t GetPoolInfo(hsa_amd_memory_pool_info_t attribute,
void* value) const;
hsa_status_t GetAgentPoolInfo(const core::Agent& agent,
hsa_amd_agent_memory_pool_info_t attribute,
void* value) const;
hsa_status_t AllowAccess(uint32_t num_agents, const hsa_agent_t* agents,
const void* ptr, size_t size) const;
hsa_status_t CanMigrate(const MemoryRegion& dst, bool& result) const;
hsa_status_t Migrate(uint32_t flag, const void* ptr) const;
hsa_status_t Lock(uint32_t num_agents, const hsa_agent_t* agents,
void* host_ptr, size_t size, void** agent_ptr) const;
hsa_status_t Unlock(void* host_ptr) const;
HSAuint64 GetBaseAddress() const { return mem_props_.VirtualBaseAddress; }
HSAuint64 GetPhysicalSize() const { return mem_props_.SizeInBytes; }
HSAuint64 GetVirtualSize() const { return virtual_size_; }
hsa_status_t AssignAgent(void* ptr, size_t size, const core::Agent& agent,
hsa_access_permission_t access) const;
__forceinline bool IsLocalMemory() const {
return ((mem_props_.HeapType == HSA_HEAPTYPE_FRAME_BUFFER_PRIVATE) ||
(mem_props_.HeapType == HSA_HEAPTYPE_FRAME_BUFFER_PUBLIC));
}
__forceinline bool IsPublic() const {
return (mem_props_.HeapType == HSA_HEAPTYPE_FRAME_BUFFER_PUBLIC);
}
__forceinline bool IsSystem() const {
return mem_props_.HeapType == HSA_HEAPTYPE_SYSTEM;
}
__forceinline bool IsLDS() const {
return mem_props_.HeapType == HSA_HEAPTYPE_GPU_LDS;
}
__forceinline bool IsGDS() const {
return mem_props_.HeapType == HSA_HEAPTYPE_GPU_GDS;
}
__forceinline bool IsScratch() const {
return mem_props_.HeapType == HSA_HEAPTYPE_GPU_SCRATCH;
}
__forceinline bool IsSvm() const {
return mem_props_.HeapType == HSA_HEAPTYPE_DEVICE_SVM;
}
__forceinline uint32_t BusWidth() const {
return static_cast<uint32_t>(mem_props_.Width);
}
__forceinline uint32_t MaxMemCloc() const {
return static_cast<uint32_t>(mem_props_.MemoryClockMax);
}
private:
const HsaMemoryProperties mem_props_;
HsaMemFlags mem_flag_;
HsaMemMapFlags map_flag_;
size_t max_single_alloc_size_;
HSAuint64 virtual_size_;
static const size_t kPageSize_ = 4096;
};
} // namespace
#endif // header guard
@@ -0,0 +1,56 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef HSA_RUNTIME_CORE_INC_AMD_TOPOLOGY_H_
#define HSA_RUNTIME_CORE_INC_AMD_TOPOLOGY_H_
namespace amd {
/// @brief Initializes the runtime.
/// Should not be called directly, must be called only from Runtime::Acquire()
bool Load();
/// @brief Shutdown/cleanup of runtime.
/// Should not be called directly, must be called only from Runtime::Release()
bool Unload();
} // namespace
#endif // header guard
+108
View File
@@ -0,0 +1,108 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef HSA_RUNTIME_CORE_INC_BLIT_H_
#define HSA_RUNTIME_CORE_INC_BLIT_H_
#include <stdint.h>
#include "core/inc/agent.h"
namespace core {
class Blit {
public:
explicit Blit() {}
virtual ~Blit() {}
/// @brief Initialize a blit object.
///
/// @param agent Pointer to the agent that will execute the blit commands.
///
/// @return hsa_status_t
virtual hsa_status_t Initialize(const core::Agent& agent) = 0;
/// @brief Marks the blit object as invalid and uncouples its link with
/// the underlying compute device's control block. Use of blit object
/// once it has been release is illegal and any behavior is indeterminate
///
/// @note: The call will block until all commands have executed.
///
/// @return hsa_status_t
virtual hsa_status_t Destroy() = 0;
/// @brief Submit a linear copy command to the the underlying compute device's
/// control block. The call is blocking until the command execution is
/// finished.
///
/// @param dst Memory address of the copy destination.
/// @param src Memory address of the copy source.
/// @param size Size of the data to be copied.
virtual hsa_status_t SubmitLinearCopyCommand(void* dst, const void* src,
size_t size) = 0;
/// @brief Submit a linear copy command to the the underlying compute device's
/// control block. The call is non blocking. The memory transfer will start
/// after all dependent signals are satisfied. After the transfer is
/// completed, the out signal will be decremented.
///
/// @param dst Memory address of the copy destination.
/// @param src Memory address of the copy source.
/// @param size Size of the data to be copied.
/// @param dep_signals Arrays of dependent signal.
/// @param out_signal Output signal.
virtual hsa_status_t SubmitLinearCopyCommand(
void* dst, const void* src, size_t size,
std::vector<core::Signal*>& dep_signals, core::Signal& out_signal) = 0;
/// @brief Submit a linear fill command to the the underlying compute device's
/// control block. The call is blocking until the command execution is
/// finished.
///
/// @param ptr Memory address of the fill destination.
/// @param value Value to be set.
/// @param num Number of uint32_t element to be set to the value.
virtual hsa_status_t SubmitLinearFillCommand(void* ptr, uint32_t value,
size_t num) = 0;
};
} // namespace core
#endif // header guard
+75
View File
@@ -0,0 +1,75 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef HSA_RUNTME_CORE_INC_CHECKED_H_
#define HSA_RUNTME_CORE_INC_CHECKED_H_
#include "stdint.h"
namespace core {
/// @brief Base class for all classes whose validity can be checked using
/// IsValid() method.
template <uint64_t code>
class Checked {
public:
typedef Checked<code> CheckedType;
Checked() { object_ = uintptr_t(this) ^ uintptr_t(code); }
Checked(const Checked&) { object_ = uintptr_t(this) ^ uintptr_t(code); }
Checked(Checked&&) { object_ = uintptr_t(this) ^ uintptr_t(code); }
virtual ~Checked() { object_ = NULL; }
const Checked& operator=(Checked&& rhs) { return *this; }
const Checked& operator=(const Checked& rhs) { return *this; }
bool IsValid() const {
return object_ == (uintptr_t(this) ^ uintptr_t(code));
}
private:
uintptr_t object_;
};
} // namespace core
#endif // header guard
@@ -0,0 +1,174 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
// HSA runtime C++ interface file.
#ifndef HSA_RUNTME_CORE_INC_DEFAULT_SIGNAL_H_
#define HSA_RUNTME_CORE_INC_DEFAULT_SIGNAL_H_
#include "core/inc/runtime.h"
#include "core/inc/signal.h"
#include "core/util/utils.h"
namespace core {
/// @brief Simple pure memory based signal.
/// @brief See base class Signal.
class DefaultSignal : public Signal {
public:
/// @brief Determines if a Signal* can be safely converted to DefaultSignal*
/// via static_cast.
static __forceinline bool IsType(Signal* ptr) {
return ptr->IsType(&rtti_id_);
}
/// @brief See base class Signal.
explicit DefaultSignal(hsa_signal_value_t initial_value);
/// @brief See base class Signal.
~DefaultSignal();
// Below are various methods corresponding to the APIs, which load/store the
// signal value or modify the existing signal value automically and with
// specified memory ordering semantics.
hsa_signal_value_t LoadRelaxed();
hsa_signal_value_t LoadAcquire();
void StoreRelaxed(hsa_signal_value_t value);
void StoreRelease(hsa_signal_value_t value);
hsa_signal_value_t WaitRelaxed(hsa_signal_condition_t condition,
hsa_signal_value_t compare_value,
uint64_t timeout, hsa_wait_state_t wait_hint);
hsa_signal_value_t WaitAcquire(hsa_signal_condition_t condition,
hsa_signal_value_t compare_value,
uint64_t timeout, hsa_wait_state_t wait_hint);
void AndRelaxed(hsa_signal_value_t value);
void AndAcquire(hsa_signal_value_t value);
void AndRelease(hsa_signal_value_t value);
void AndAcqRel(hsa_signal_value_t value);
void OrRelaxed(hsa_signal_value_t value);
void OrAcquire(hsa_signal_value_t value);
void OrRelease(hsa_signal_value_t value);
void OrAcqRel(hsa_signal_value_t value);
void XorRelaxed(hsa_signal_value_t value);
void XorAcquire(hsa_signal_value_t value);
void XorRelease(hsa_signal_value_t value);
void XorAcqRel(hsa_signal_value_t value);
void AddRelaxed(hsa_signal_value_t value);
void AddAcquire(hsa_signal_value_t value);
void AddRelease(hsa_signal_value_t value);
void AddAcqRel(hsa_signal_value_t value);
void SubRelaxed(hsa_signal_value_t value);
void SubAcquire(hsa_signal_value_t value);
void SubRelease(hsa_signal_value_t value);
void SubAcqRel(hsa_signal_value_t value);
hsa_signal_value_t ExchRelaxed(hsa_signal_value_t value);
hsa_signal_value_t ExchAcquire(hsa_signal_value_t value);
hsa_signal_value_t ExchRelease(hsa_signal_value_t value);
hsa_signal_value_t ExchAcqRel(hsa_signal_value_t value);
hsa_signal_value_t CasRelaxed(hsa_signal_value_t expected,
hsa_signal_value_t value);
hsa_signal_value_t CasAcquire(hsa_signal_value_t expected,
hsa_signal_value_t value);
hsa_signal_value_t CasRelease(hsa_signal_value_t expected,
hsa_signal_value_t value);
hsa_signal_value_t CasAcqRel(hsa_signal_value_t expected,
hsa_signal_value_t value);
/// @brief see the base class Signal
__forceinline hsa_signal_value_t* ValueLocation() const {
return (hsa_signal_value_t*)&signal_.value;
}
/// @brief see the base class Signal
__forceinline HsaEvent* EopEvent() { return NULL; }
/// @brief prevent throwing exceptions
void* operator new(size_t size) { return malloc(size); }
/// @brief prevent throwing exceptions
void operator delete(void* ptr) { free(ptr); }
protected:
bool _IsA(rtti_t id) const { return id == &rtti_id_; }
private:
static int rtti_id_;
DISALLOW_COPY_AND_ASSIGN(DefaultSignal);
};
} // namespace core
#endif // header guard
+167
View File
@@ -0,0 +1,167 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef HSA_RUNTIME_CORE_INC_HOST_QUEUE_H_
#define HSA_RUNTIME_CORE_INC_HOST_QUEUE_H_
#include "core/inc/memory_region.h"
#include "core/inc/queue.h"
#include "core/inc/runtime.h"
#include "core/inc/signal.h"
namespace core {
class HostQueue : public Queue {
public:
HostQueue(hsa_region_t region, uint32_t ring_size, hsa_queue_type_t type,
uint32_t features, hsa_signal_t doorbell_signal);
~HostQueue();
hsa_status_t Inactivate() { return HSA_STATUS_SUCCESS; }
uint64_t LoadReadIndexAcquire() {
return atomic::Load(&amd_queue_.read_dispatch_id,
std::memory_order_acquire);
}
uint64_t LoadReadIndexRelaxed() {
return atomic::Load(&amd_queue_.read_dispatch_id,
std::memory_order_relaxed);
}
uint64_t LoadWriteIndexAcquire() {
return atomic::Load(&amd_queue_.write_dispatch_id,
std::memory_order_acquire);
}
uint64_t LoadWriteIndexRelaxed() {
return atomic::Load(&amd_queue_.write_dispatch_id,
std::memory_order_relaxed);
}
void StoreReadIndexRelaxed(uint64_t value) {
atomic::Store(&amd_queue_.read_dispatch_id, value,
std::memory_order_relaxed);
}
void StoreReadIndexRelease(uint64_t value) {
atomic::Store(&amd_queue_.read_dispatch_id, value,
std::memory_order_release);
}
void StoreWriteIndexRelaxed(uint64_t value) {
atomic::Store(&amd_queue_.write_dispatch_id, value,
std::memory_order_relaxed);
}
void StoreWriteIndexRelease(uint64_t value) {
atomic::Store(&amd_queue_.write_dispatch_id, value,
std::memory_order_release);
}
uint64_t CasWriteIndexAcqRel(uint64_t expected, uint64_t value) {
return atomic::Cas(&amd_queue_.write_dispatch_id, value, expected,
std::memory_order_acq_rel);
}
uint64_t CasWriteIndexAcquire(uint64_t expected, uint64_t value) {
return atomic::Cas(&amd_queue_.write_dispatch_id, value, expected,
std::memory_order_acquire);
}
uint64_t CasWriteIndexRelaxed(uint64_t expected, uint64_t value) {
return atomic::Cas(&amd_queue_.write_dispatch_id, value, expected,
std::memory_order_relaxed);
}
uint64_t CasWriteIndexRelease(uint64_t expected, uint64_t value) {
return atomic::Cas(&amd_queue_.write_dispatch_id, value, expected,
std::memory_order_release);
}
uint64_t AddWriteIndexAcqRel(uint64_t value) {
return atomic::Add(&amd_queue_.write_dispatch_id, value,
std::memory_order_acq_rel);
}
uint64_t AddWriteIndexAcquire(uint64_t value) {
return atomic::Add(&amd_queue_.write_dispatch_id, value,
std::memory_order_acquire);
}
uint64_t AddWriteIndexRelaxed(uint64_t value) {
return atomic::Add(&amd_queue_.write_dispatch_id, value,
std::memory_order_relaxed);
}
uint64_t AddWriteIndexRelease(uint64_t value) {
return atomic::Add(&amd_queue_.write_dispatch_id, value,
std::memory_order_release);
}
hsa_status_t SetCUMasking(const uint32_t num_cu_mask_count,
const uint32_t* cu_mask) {
return HSA_STATUS_ERROR;
}
bool active() const { return active_; }
void* operator new(size_t size) {
return _aligned_malloc(size, HSA_QUEUE_ALIGN_BYTES);
}
void* operator new(size_t size, void* ptr) { return ptr; }
void operator delete(void* ptr) { _aligned_free(ptr); }
void operator delete(void*, void*) {}
private:
static const size_t kRingAlignment = 256;
const uint32_t size_;
bool active_;
void* ring_;
DISALLOW_COPY_AND_ASSIGN(HostQueue);
};
} // namespace core
#endif // header guard
@@ -0,0 +1,63 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef HSA_RUNTIME_CORE_INC_HSA_API_TRACE_INT_H
#define HSA_RUNTIME_CORE_INC_HSA_API_TRACE_INT_H
#include "inc/hsa_api_trace.h"
#include "core/inc/hsa_internal.h"
namespace core {
struct ApiTable {
::ApiTable table;
ExtTable extension_backup;
ApiTable();
void Reset();
void LinkExts(ExtTable* ptr);
};
extern ApiTable hsa_api_table_;
extern ApiTable hsa_internal_api_table_;
}
#endif
@@ -0,0 +1,80 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef HSA_RUNTME_CORE_INC_AMD_EXT_INTERFACE_H_
#define HSA_RUNTME_CORE_INC_AMD_EXT_INTERFACE_H_
#include <string>
#include <vector>
#include "hsa_api_trace_int.h"
#include "core/util/os.h"
#include "core/util/utils.h"
namespace core {
struct ExtTableInternal : public ExtTable {
decltype(::hsa_amd_image_get_info_max_dim)* hsa_amd_image_get_info_max_dim_fn;
decltype(::hsa_amd_image_create)* hsa_amd_image_create_fn;
};
class ExtensionEntryPoints {
public:
ExtTableInternal table;
ExtensionEntryPoints();
bool Load(std::string library_name);
void Unload();
private:
typedef void (*Load_t)(const ::ApiTable* table);
typedef void (*Unload_t)();
std::vector<os::LibHandle> libs_;
void InitTable();
DISALLOW_COPY_AND_ASSIGN(ExtensionEntryPoints);
};
}
#endif
+347
View File
@@ -0,0 +1,347 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef HSA_RUNTIME_CORE_INC_HSA_INTERNAL_H
#define HSA_RUNTIME_CORE_INC_HSA_INTERNAL_H
#include "inc/hsa.h"
namespace HSA
{
// Define core namespace interfaces - copy of function declarations in hsa.h
hsa_status_t HSA_API hsa_init();
hsa_status_t HSA_API hsa_shut_down();
hsa_status_t HSA_API
hsa_system_get_info(hsa_system_info_t attribute, void *value);
hsa_status_t HSA_API
hsa_system_extension_supported(uint16_t extension, uint16_t version_major,
uint16_t version_minor, bool *result);
hsa_status_t HSA_API
hsa_system_get_extension_table(uint16_t extension, uint16_t version_major,
uint16_t version_minor, void *table);
hsa_status_t HSA_API
hsa_iterate_agents(hsa_status_t (*callback)(hsa_agent_t agent, void *data),
void *data);
hsa_status_t HSA_API hsa_agent_get_info(hsa_agent_t agent,
hsa_agent_info_t attribute,
void *value);
hsa_status_t HSA_API hsa_agent_get_exception_policies(hsa_agent_t agent,
hsa_profile_t profile,
uint16_t *mask);
hsa_status_t HSA_API
hsa_agent_extension_supported(uint16_t extension, hsa_agent_t agent,
uint16_t version_major,
uint16_t version_minor, bool *result);
hsa_status_t HSA_API
hsa_queue_create(hsa_agent_t agent, uint32_t size, hsa_queue_type_t type,
void (*callback)(hsa_status_t status, hsa_queue_t *source,
void *data),
void *data, uint32_t private_segment_size,
uint32_t group_segment_size, hsa_queue_t **queue);
hsa_status_t HSA_API
hsa_soft_queue_create(hsa_region_t region, uint32_t size,
hsa_queue_type_t type, uint32_t features,
hsa_signal_t completion_signal, hsa_queue_t **queue);
hsa_status_t HSA_API hsa_queue_destroy(hsa_queue_t *queue);
hsa_status_t HSA_API hsa_queue_inactivate(hsa_queue_t *queue);
uint64_t HSA_API hsa_queue_load_read_index_acquire(const hsa_queue_t *queue);
uint64_t HSA_API hsa_queue_load_read_index_relaxed(const hsa_queue_t *queue);
uint64_t HSA_API hsa_queue_load_write_index_acquire(const hsa_queue_t *queue);
uint64_t HSA_API hsa_queue_load_write_index_relaxed(const hsa_queue_t *queue);
void HSA_API hsa_queue_store_write_index_relaxed(const hsa_queue_t *queue,
uint64_t value);
void HSA_API hsa_queue_store_write_index_release(const hsa_queue_t *queue,
uint64_t value);
uint64_t HSA_API hsa_queue_cas_write_index_acq_rel(const hsa_queue_t *queue,
uint64_t expected,
uint64_t value);
uint64_t HSA_API hsa_queue_cas_write_index_acquire(const hsa_queue_t *queue,
uint64_t expected,
uint64_t value);
uint64_t HSA_API hsa_queue_cas_write_index_relaxed(const hsa_queue_t *queue,
uint64_t expected,
uint64_t value);
uint64_t HSA_API hsa_queue_cas_write_index_release(const hsa_queue_t *queue,
uint64_t expected,
uint64_t value);
uint64_t HSA_API
hsa_queue_add_write_index_acq_rel(const hsa_queue_t *queue, uint64_t value);
uint64_t HSA_API
hsa_queue_add_write_index_acquire(const hsa_queue_t *queue, uint64_t value);
uint64_t HSA_API
hsa_queue_add_write_index_relaxed(const hsa_queue_t *queue, uint64_t value);
uint64_t HSA_API
hsa_queue_add_write_index_release(const hsa_queue_t *queue, uint64_t value);
void HSA_API hsa_queue_store_read_index_relaxed(const hsa_queue_t *queue,
uint64_t value);
void HSA_API hsa_queue_store_read_index_release(const hsa_queue_t *queue,
uint64_t value);
hsa_status_t HSA_API hsa_agent_iterate_regions(
hsa_agent_t agent,
hsa_status_t (*callback)(hsa_region_t region, void *data), void *data);
hsa_status_t HSA_API hsa_region_get_info(hsa_region_t region,
hsa_region_info_t attribute,
void *value);
hsa_status_t HSA_API hsa_memory_register(void *address, size_t size);
hsa_status_t HSA_API hsa_memory_deregister(void *address, size_t size);
hsa_status_t HSA_API
hsa_memory_allocate(hsa_region_t region, size_t size, void **ptr);
hsa_status_t HSA_API hsa_memory_free(void *ptr);
hsa_status_t HSA_API hsa_memory_copy(void *dst, const void *src, size_t size);
hsa_status_t HSA_API hsa_memory_assign_agent(void *ptr, hsa_agent_t agent,
hsa_access_permission_t access);
hsa_status_t HSA_API
hsa_signal_create(hsa_signal_value_t initial_value, uint32_t num_consumers,
const hsa_agent_t *consumers, hsa_signal_t *signal);
hsa_status_t HSA_API hsa_signal_destroy(hsa_signal_t signal);
hsa_signal_value_t HSA_API hsa_signal_load_relaxed(hsa_signal_t signal);
hsa_signal_value_t HSA_API hsa_signal_load_acquire(hsa_signal_t signal);
void HSA_API
hsa_signal_store_relaxed(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_store_release(hsa_signal_t signal, hsa_signal_value_t value);
hsa_signal_value_t HSA_API
hsa_signal_wait_relaxed(hsa_signal_t signal,
hsa_signal_condition_t condition,
hsa_signal_value_t compare_value,
uint64_t timeout_hint,
hsa_wait_state_t wait_expectancy_hint);
hsa_signal_value_t HSA_API
hsa_signal_wait_acquire(hsa_signal_t signal,
hsa_signal_condition_t condition,
hsa_signal_value_t compare_value,
uint64_t timeout_hint,
hsa_wait_state_t wait_expectancy_hint);
void HSA_API
hsa_signal_and_relaxed(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_and_acquire(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_and_release(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_and_acq_rel(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_or_relaxed(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_or_acquire(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_or_release(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_or_acq_rel(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_xor_relaxed(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_xor_acquire(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_xor_release(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_xor_acq_rel(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_add_relaxed(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_add_acquire(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_add_release(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_add_acq_rel(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_subtract_relaxed(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_subtract_acquire(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_subtract_release(hsa_signal_t signal, hsa_signal_value_t value);
void HSA_API
hsa_signal_subtract_acq_rel(hsa_signal_t signal, hsa_signal_value_t value);
hsa_signal_value_t HSA_API
hsa_signal_exchange_relaxed(hsa_signal_t signal, hsa_signal_value_t value);
hsa_signal_value_t HSA_API
hsa_signal_exchange_acquire(hsa_signal_t signal, hsa_signal_value_t value);
hsa_signal_value_t HSA_API
hsa_signal_exchange_release(hsa_signal_t signal, hsa_signal_value_t value);
hsa_signal_value_t HSA_API
hsa_signal_exchange_acq_rel(hsa_signal_t signal, hsa_signal_value_t value);
hsa_signal_value_t HSA_API hsa_signal_cas_relaxed(hsa_signal_t signal,
hsa_signal_value_t expected,
hsa_signal_value_t value);
hsa_signal_value_t HSA_API hsa_signal_cas_acquire(hsa_signal_t signal,
hsa_signal_value_t expected,
hsa_signal_value_t value);
hsa_signal_value_t HSA_API hsa_signal_cas_release(hsa_signal_t signal,
hsa_signal_value_t expected,
hsa_signal_value_t value);
hsa_signal_value_t HSA_API hsa_signal_cas_acq_rel(hsa_signal_t signal,
hsa_signal_value_t expected,
hsa_signal_value_t value);
hsa_status_t hsa_isa_from_name(
const char *name,
hsa_isa_t *isa
);
hsa_status_t HSA_API hsa_isa_get_info(
hsa_isa_t isa,
hsa_isa_info_t attribute,
uint32_t index,
void *value
);
hsa_status_t hsa_isa_compatible(
hsa_isa_t code_object_isa,
hsa_isa_t agent_isa,
bool *result
);
hsa_status_t HSA_API hsa_code_object_serialize(
hsa_code_object_t code_object,
hsa_status_t (*alloc_callback)(
size_t size, hsa_callback_data_t data, void **address
),
hsa_callback_data_t callback_data,
const char *options,
void **serialized_code_object,
size_t *serialized_code_object_size
);
hsa_status_t HSA_API hsa_code_object_deserialize(
void *serialized_code_object,
size_t serialized_code_object_size,
const char *options,
hsa_code_object_t *code_object
);
hsa_status_t HSA_API hsa_code_object_destroy(
hsa_code_object_t code_object
);
hsa_status_t HSA_API hsa_code_object_get_info(
hsa_code_object_t code_object,
hsa_code_object_info_t attribute,
void *value
);
hsa_status_t HSA_API hsa_code_object_get_symbol(
hsa_code_object_t code_object,
const char *symbol_name,
hsa_code_symbol_t *symbol
);
hsa_status_t HSA_API hsa_code_symbol_get_info(
hsa_code_symbol_t code_symbol,
hsa_code_symbol_info_t attribute,
void *value
);
hsa_status_t HSA_API hsa_code_object_iterate_symbols(
hsa_code_object_t code_object,
hsa_status_t (*callback)(
hsa_code_object_t code_object, hsa_code_symbol_t symbol, void *data
),
void *data
);
hsa_status_t HSA_API hsa_executable_create(
hsa_profile_t profile,
hsa_executable_state_t executable_state,
const char *options,
hsa_executable_t *executable
);
hsa_status_t HSA_API hsa_executable_destroy(
hsa_executable_t executable
);
hsa_status_t HSA_API hsa_executable_load_code_object(
hsa_executable_t executable,
hsa_agent_t agent,
hsa_code_object_t code_object,
const char *options
);
hsa_status_t HSA_API hsa_executable_freeze(
hsa_executable_t executable,
const char *options
);
hsa_status_t HSA_API hsa_executable_get_info(
hsa_executable_t executable,
hsa_executable_info_t attribute,
void *value
);
hsa_status_t HSA_API hsa_executable_global_variable_define(
hsa_executable_t executable,
const char *variable_name,
void *address
);
hsa_status_t HSA_API hsa_executable_agent_global_variable_define(
hsa_executable_t executable,
hsa_agent_t agent,
const char *variable_name,
void *address
);
hsa_status_t HSA_API hsa_executable_readonly_variable_define(
hsa_executable_t executable,
hsa_agent_t agent,
const char *variable_name,
void *address
);
hsa_status_t HSA_API hsa_executable_validate(
hsa_executable_t executable,
uint32_t *result
);
hsa_status_t HSA_API hsa_executable_get_symbol(
hsa_executable_t executable,
const char *module_name,
const char *symbol_name,
hsa_agent_t agent,
int32_t call_convention,
hsa_executable_symbol_t *symbol
);
hsa_status_t HSA_API hsa_executable_symbol_get_info(
hsa_executable_symbol_t executable_symbol,
hsa_executable_symbol_info_t attribute,
void *value
);
hsa_status_t HSA_API hsa_executable_iterate_symbols(
hsa_executable_t executable,
hsa_status_t (*callback)(
hsa_executable_t executable, hsa_executable_symbol_t symbol, void *data
),
void *data
);
hsa_status_t HSA_API
hsa_status_string(hsa_status_t status, const char **status_string);
}
#ifdef BUILDING_HSA_CORE_RUNTIME
//This using declaration is deliberate!
//We want unqualified name resolution to fail when building the runtime. This is a guard against accidental use of the intercept layer in the runtime.
using namespace HSA;
#endif
#endif
@@ -0,0 +1,47 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#include "hsa_api_trace.h"
void hsa_table_interface_init(const ApiTable* table);
const ApiTable* hsa_table_interface_get_table();
@@ -0,0 +1,206 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
// HSA runtime C++ interface file.
#ifndef HSA_RUNTME_CORE_INC_INTERRUPT_SIGNAL_H_
#define HSA_RUNTME_CORE_INC_INTERRUPT_SIGNAL_H_
#include "hsakmt.h"
#include "core/inc/runtime.h"
#include "core/inc/signal.h"
#include "core/util/utils.h"
namespace core {
/// @brief A Signal implementation using interrupts versus plain memory based.
/// Also see base class Signal.
///
/// Breaks common/vendor separation - signals in general needs to be re-worked
/// at the foundation level to make sense in a multi-device system.
/// Supports only one waiter for now.
/// KFD changes are needed to support multiple waiters and have device
/// signaling.
class InterruptSignal : public Signal {
public:
static HsaEvent* CreateEvent(HSA_EVENTTYPE type, bool manual_reset);
static void DestroyEvent(HsaEvent* evt);
/// @brief Determines if a Signal* can be safely converted to an
/// InterruptSignal* via static_cast.
static __forceinline bool IsType(Signal* ptr) {
return ptr->IsType(&rtti_id_);
}
explicit InterruptSignal(hsa_signal_value_t initial_value,
HsaEvent* use_event = NULL);
~InterruptSignal();
// Below are various methods corresponding to the APIs, which load/store the
// signal value or modify the existing signal value automically and with
// specified memory ordering semantics.
hsa_signal_value_t LoadRelaxed();
hsa_signal_value_t LoadAcquire();
void StoreRelaxed(hsa_signal_value_t value);
void StoreRelease(hsa_signal_value_t value);
hsa_signal_value_t WaitRelaxed(hsa_signal_condition_t condition,
hsa_signal_value_t compare_value,
uint64_t timeout, hsa_wait_state_t wait_hint);
hsa_signal_value_t WaitAcquire(hsa_signal_condition_t condition,
hsa_signal_value_t compare_value,
uint64_t timeout, hsa_wait_state_t wait_hint);
void AndRelaxed(hsa_signal_value_t value);
void AndAcquire(hsa_signal_value_t value);
void AndRelease(hsa_signal_value_t value);
void AndAcqRel(hsa_signal_value_t value);
void OrRelaxed(hsa_signal_value_t value);
void OrAcquire(hsa_signal_value_t value);
void OrRelease(hsa_signal_value_t value);
void OrAcqRel(hsa_signal_value_t value);
void XorRelaxed(hsa_signal_value_t value);
void XorAcquire(hsa_signal_value_t value);
void XorRelease(hsa_signal_value_t value);
void XorAcqRel(hsa_signal_value_t value);
void AddRelaxed(hsa_signal_value_t value);
void AddAcquire(hsa_signal_value_t value);
void AddRelease(hsa_signal_value_t value);
void AddAcqRel(hsa_signal_value_t value);
void SubRelaxed(hsa_signal_value_t value);
void SubAcquire(hsa_signal_value_t value);
void SubRelease(hsa_signal_value_t value);
void SubAcqRel(hsa_signal_value_t value);
hsa_signal_value_t ExchRelaxed(hsa_signal_value_t value);
hsa_signal_value_t ExchAcquire(hsa_signal_value_t value);
hsa_signal_value_t ExchRelease(hsa_signal_value_t value);
hsa_signal_value_t ExchAcqRel(hsa_signal_value_t value);
hsa_signal_value_t CasRelaxed(hsa_signal_value_t expected,
hsa_signal_value_t value);
hsa_signal_value_t CasAcquire(hsa_signal_value_t expected,
hsa_signal_value_t value);
hsa_signal_value_t CasRelease(hsa_signal_value_t expected,
hsa_signal_value_t value);
hsa_signal_value_t CasAcqRel(hsa_signal_value_t expected,
hsa_signal_value_t value);
/// @brief See base class Signal.
__forceinline hsa_signal_value_t* ValueLocation() const {
return (hsa_signal_value_t*)&signal_.value;
}
/// @brief See base class Signal.
__forceinline HsaEvent* EopEvent() { return event_; }
// TODO(bwicakso) : work around for SDMA async copy. Bypass waiting on EOP
// event because SDMA copy does not handle interrupt yet.
__forceinline void DisableWaitEvent() { wait_on_event_ = false; }
/// @brief prevent throwing exceptions
void* operator new(size_t size) { return malloc(size); }
/// @brief prevent throwing exceptions
void operator delete(void* ptr) { free(ptr); }
protected:
bool _IsA(rtti_t id) const { return id == &rtti_id_; }
private:
/// @variable KFD event on which the interrupt signal is based on.
HsaEvent* event_;
/// @variable Indicates whether the signal should release the event when it
/// closes or not.
bool free_event_;
// TODO(bwicakso) : work around for SDMA async copy. Bypass waiting on EOP
// event because SDMA copy does not handle interrupt yet.
bool wait_on_event_;
/// Used to obtain a globally unique value (address) for rtti.
static int rtti_id_;
/// @brief Notify driver of signal value change if necessary.
__forceinline void SetEvent() {
std::atomic_signal_fence(std::memory_order_seq_cst);
if (InWaiting()) hsaKmtSetEvent(event_);
}
DISALLOW_COPY_AND_ASSIGN(InterruptSignal);
};
} // namespace core
#endif // header guard
+164
View File
@@ -0,0 +1,164 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
#ifndef HSA_RUNTIME_CORE_ISA_H_
#define HSA_RUNTIME_CORE_ISA_H_
#include <cassert>
#include <cstdint>
#include <string>
#include <tuple>
#include <unordered_map>
#include "core/inc/amd_hsa_code.hpp"
namespace core {
// @class Isa
// @brief Instruction Set Architecture
class Isa final: public amd::hsa::common::Signed<0xB13594F2BD8F212D> {
public:
// @brief Isa's version type
typedef std::tuple<int32_t, int32_t, int32_t> Version;
// @brief Default destructor
~Isa() {}
// @returns Handle equivalent of @p isa_object
static hsa_isa_t Handle(const Isa *isa_object) {
hsa_isa_t isa_handle = { reinterpret_cast<uint64_t>(isa_object) };
return isa_handle;
}
// @returns Object equivalend of @p isa_handle
static Isa *Object(const hsa_isa_t &isa_handle) {
Isa *isa_object = amd::hsa::common::ObjectAt<Isa>(isa_handle.handle);
return isa_object;
}
// @returns This Isa's version
const Version &version() const {
return version_;
}
// @returns This Isa's vendor
std::string GetVendor() const {
return "AMD";
}
// @returns This Isa's architecture
std::string GetArchitecture() const {
return "AMDGPU";
}
// @returns This Isa's major version
int32_t GetMajorVersion() const {
return std::get<0>(version_);
}
// @returns This Isa's minor version
int32_t GetMinorVersion() const {
return std::get<1>(version_);
}
// @returns This Isa's stepping
int32_t GetStepping() const {
return std::get<2>(version_);
}
// @returns True if this Isa is compatible with @p isa_object, false otherwise
bool IsCompatible(const Isa *isa_object) const {
assert(isa_object);
return version_ == isa_object->version_;
}
// @returns True if this Isa is compatible with @p isa_handle, false otherwise
bool IsCompatible(const hsa_isa_t &isa_handle) const {
assert(isa_handle.handle);
return IsCompatible(Object(isa_handle));
}
// @brief Isa is always in valid state
bool IsValid() const {
return true;
}
// @returns This Isa's full name
std::string GetFullName() const;
// @brief Query value of requested @p attribute and record it in @p value
bool GetInfo(const hsa_isa_info_t &attribute, void *value) const;
private:
// @brief Default constructor
Isa(): version_(Version(-1, -1, -1)) {}
// @brief Construct from @p version
Isa(const Version &version): version_(version) {}
// @brief Isa's version
Version version_;
// @brief Isa's friends
friend class IsaRegistry;
}; // class Isa
// @class IsaRegistry
// @brief Instruction Set Architecture Registry
class IsaRegistry final {
public:
// @returns Isa for requested @p full_name, null pointer if not supported
static const Isa *GetIsa(const std::string &full_name);
// @returns Isa for requested @p version, null pointer if not supported
static const Isa *GetIsa(const Isa::Version &version);
private:
// @brief IsaRegistry's map type
typedef std::unordered_map<std::string, Isa> IsaMap;
// @brief Supported instruction set architectures
static const IsaMap supported_isas_;
// @brief Default constructor - not available
IsaRegistry();
// @brief Default destructor - not available
~IsaRegistry();
// @returns Supported instruction set architectures
static const IsaMap GetSupportedIsas();
}; // class IsaRegistry
} // namespace core
#endif // HSA_RUNTIME_CORE_ISA_HPP_
@@ -0,0 +1,109 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
// HSA runtime C++ interface file.
#ifndef HSA_RUNTME_CORE_INC_MEMORY_REGION_H_
#define HSA_RUNTME_CORE_INC_MEMORY_REGION_H_
#include <vector>
#include "core/inc/runtime.h"
#include "core/inc/agent.h"
#include "core/inc/checked.h"
namespace core {
class Agent;
class MemoryRegion : public Checked<0x9C961F19EE175BB3> {
public:
MemoryRegion(bool fine_grain, bool full_profile, core::Agent* owner)
: fine_grain_(fine_grain), full_profile_(full_profile), owner_(owner) {
assert(owner_ != NULL);
}
virtual ~MemoryRegion() {}
// Convert this object into hsa_region_t.
static __forceinline hsa_region_t Convert(MemoryRegion* region) {
const hsa_region_t region_handle = {
static_cast<uint64_t>(reinterpret_cast<uintptr_t>(region))};
return region_handle;
}
static __forceinline const hsa_region_t Convert(const MemoryRegion* region) {
const hsa_region_t region_handle = {
static_cast<uint64_t>(reinterpret_cast<uintptr_t>(region))};
return region_handle;
}
// Convert hsa_region_t into MemoryRegion *.
static __forceinline MemoryRegion* Convert(hsa_region_t region) {
return reinterpret_cast<MemoryRegion*>(region.handle);
}
virtual hsa_status_t Allocate(size_t size, void** address) const = 0;
virtual hsa_status_t Free(void* address, size_t size) const = 0;
// Translate memory properties into HSA region attribute.
virtual hsa_status_t GetInfo(hsa_region_info_t attribute,
void* value) const = 0;
virtual hsa_status_t AssignAgent(void* ptr, size_t size, const Agent& agent,
hsa_access_permission_t access) const = 0;
__forceinline bool fine_grain() const { return fine_grain_; }
__forceinline bool full_profile() const { return full_profile_; }
__forceinline core::Agent* owner() const { return owner_; }
private:
const bool fine_grain_;
const bool full_profile_;
core::Agent* owner_;
};
} // namespace core
#endif // header guard
+322
View File
@@ -0,0 +1,322 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
// HSA runtime C++ interface file.
#ifndef HSA_RUNTME_CORE_INC_COMMAND_QUEUE_H_
#define HSA_RUNTME_CORE_INC_COMMAND_QUEUE_H_
#include <sstream>
#include "core/common/shared.h"
#include "core/inc/runtime.h"
#include "core/inc/checked.h"
#include "core/util/utils.h"
#include "inc/amd_hsa_queue.h"
namespace core {
struct AqlPacket {
union {
hsa_kernel_dispatch_packet_t dispatch;
hsa_barrier_and_packet_t barrier_and;
hsa_barrier_or_packet_t barrier_or;
hsa_agent_dispatch_packet_t agent;
};
uint8_t type() {
return ((dispatch.header >> HSA_PACKET_HEADER_TYPE) &
((1 << HSA_PACKET_HEADER_WIDTH_TYPE) - 1));
}
bool IsValid() {
const uint8_t packet_type = dispatch.header >> HSA_PACKET_HEADER_TYPE;
return (packet_type > HSA_PACKET_TYPE_INVALID &&
packet_type <= HSA_PACKET_TYPE_BARRIER_OR);
}
std::string string() const {
std::stringstream string;
uint8_t type = ((dispatch.header >> HSA_PACKET_HEADER_TYPE) &
((1 << HSA_PACKET_HEADER_WIDTH_TYPE) - 1));
const char* type_names[] = {
"HSA_PACKET_TYPE_VENDOR_SPECIFIC", "HSA_PACKET_TYPE_INVALID",
"HSA_PACKET_TYPE_KERNEL_DISPATCH", "HSA_PACKET_TYPE_BARRIER_AND",
"HSA_PACKET_TYPE_AGENT_DISPATCH", "HSA_PACKET_TYPE_BARRIER_OR"};
string << "type: " << type_names[type]
<< "\nbarrier: " << ((dispatch.header >> HSA_PACKET_HEADER_BARRIER) &
((1 << HSA_PACKET_HEADER_WIDTH_BARRIER) - 1))
<< "\nacquire: "
<< ((dispatch.header >> HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) &
((1 << HSA_PACKET_HEADER_WIDTH_ACQUIRE_FENCE_SCOPE) - 1))
<< "\nrelease: "
<< ((dispatch.header >> HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE) &
((1 << HSA_PACKET_HEADER_WIDTH_RELEASE_FENCE_SCOPE) - 1));
if (type == HSA_PACKET_TYPE_KERNEL_DISPATCH) {
string << "\nDim: " << dispatch.setup
<< "\nworkgroup_size: " << dispatch.workgroup_size_x << ", "
<< dispatch.workgroup_size_y << ", " << dispatch.workgroup_size_z
<< "\ngrid_size: " << dispatch.grid_size_x << ", "
<< dispatch.grid_size_y << ", " << dispatch.grid_size_z
<< "\nprivate_size: " << dispatch.private_segment_size
<< "\ngroup_size: " << dispatch.group_segment_size
<< "\nkernel_object: " << dispatch.kernel_object
<< "\nkern_arg: " << dispatch.kernarg_address
<< "\nsignal: " << dispatch.completion_signal.handle;
}
if ((type == HSA_PACKET_TYPE_BARRIER_AND) ||
(type == HSA_PACKET_TYPE_BARRIER_OR)) {
for (int i = 0; i < 5; i++)
string << "\ndep[" << i << "]: " << barrier_and.dep_signal[i].handle;
string << "\nsignal: " << barrier_and.completion_signal.handle;
}
return string.str();
}
};
class Queue;
/// @brief Helper structure to simplify conversion of amd_queue_t and
/// core::Queue object.
struct SharedQueue {
amd_queue_t amd_queue;
Queue* core_queue;
};
/// @brief Class Queue which encapsulate user mode queues and
/// provides Api to access its Read, Write indices using Acquire,
/// Release and Relaxed semantics.
/*
Queue is intended to be an pure interface class and may be wrapped or replaced
by tools.
All funtions other than Convert and public_handle must be virtual.
*/
class Queue : public Checked<0xFA3906A679F9DB49>,
public Shared<SharedQueue, AMD_QUEUE_ALIGN_BYTES> {
public:
Queue() : Shared(), amd_queue_(shared_object()->amd_queue) {
if (!Shared::IsSharedObjectAllocationValid()) {
return;
}
shared_object()->core_queue = this;
public_handle_ = Convert(this);
}
virtual ~Queue() {}
/// @brief Returns the handle of Queue's public data type
///
/// @param queue Pointer to an instance of Queue implementation object
///
/// @return hsa_queue_t * Pointer to the public data type of a queue
static __forceinline hsa_queue_t* Convert(Queue* queue) {
return ((queue != NULL) && (queue->IsSharedObjectAllocationValid()))
? &queue->amd_queue_.hsa_queue
: NULL;
}
/// @brief Transform the public data type of a Queue's data type into an
// instance of it Queue class object
///
/// @param queue Handle of public data type of a queue
///
/// @return Queue * Pointer to the Queue's implementation object
static __forceinline Queue* Convert(const hsa_queue_t* queue) {
return (queue != NULL)
? reinterpret_cast<const SharedQueue*>(
reinterpret_cast<uintptr_t>(queue) -
(reinterpret_cast<uintptr_t>(
&reinterpret_cast<SharedQueue*>(1234)
->amd_queue.hsa_queue) -
uintptr_t(1234)))->core_queue
: NULL;
}
/// @brief Inactivate the queue object. Once inactivate a
/// queue cannot be used anymore and must be destroyed
///
/// @return hsa_status_t Status of request
virtual hsa_status_t Inactivate() = 0;
/// @brief Reads the Read Index of Queue using Acquire semantics
///
/// @return uint64_t Value of Read index
virtual uint64_t LoadReadIndexAcquire() = 0;
/// @brief Reads the Read Index of Queue using Relaxed semantics
///
/// @return uint64_t Value of Read index
virtual uint64_t LoadReadIndexRelaxed() = 0;
/// @brief Reads the Write Index of Queue using Acquire semantics
///
/// @return uint64_t Value of Write index
virtual uint64_t LoadWriteIndexAcquire() = 0;
/// Reads the Write Index of Queue using Relaxed semantics
///
/// @return uint64_t Value of Write index
virtual uint64_t LoadWriteIndexRelaxed() = 0;
/// @brief Updates the Read Index of Queue using Relaxed semantics
///
/// @param value New value of Read index to update
virtual void StoreReadIndexRelaxed(uint64_t value) = 0;
/// @brief Updates the Read Index of Queue using Release semantics
///
/// @param value New value of Read index to update
virtual void StoreReadIndexRelease(uint64_t value) = 0;
/// @brief Updates the Write Index of Queue using Relaxed semantics
///
/// @param value New value of Write index to update
virtual void StoreWriteIndexRelaxed(uint64_t value) = 0;
/// @brief Updates the Write Index of Queue using Release semantics
///
/// @param value New value of Write index to update
virtual void StoreWriteIndexRelease(uint64_t value) = 0;
/// @brief Compares and swaps Write index using Acquire and Release semantics
///
/// @param expected Current value of write index
///
/// @param value Value of new write index
///
/// @return uint64_t Value of write index before the update
virtual uint64_t CasWriteIndexAcqRel(uint64_t expected, uint64_t value) = 0;
/// @brief Compares and swaps Write index using Acquire semantics
///
/// @param expected Current value of write index
///
/// @param value Value of new write index
///
/// @return uint64_t Value of write index before the update
virtual uint64_t CasWriteIndexAcquire(uint64_t expected, uint64_t value) = 0;
/// @brief Compares and swaps Write index using Relaxed semantics
///
/// @param expected Current value of write index
///
/// @param value Value of new write index
///
/// @return uint64_t Value of write index before the update
virtual uint64_t CasWriteIndexRelaxed(uint64_t expected, uint64_t value) = 0;
/// @brief Compares and swaps Write index using Release semantics
///
/// @param expected Current value of write index
///
/// @param value Value of new write index
///
/// @return uint64_t Value of write index before the update
virtual uint64_t CasWriteIndexRelease(uint64_t expected, uint64_t value) = 0;
/// @brief Updates the Write index using Acquire and Release semantics
///
/// @param value Value of new write index
///
/// @return uint64_t Value of write index before the update
virtual uint64_t AddWriteIndexAcqRel(uint64_t value) = 0;
/// @brief Updates the Write index using Acquire semantics
///
/// @param value Value of new write index
///
/// @return uint64_t Value of write index before the update
virtual uint64_t AddWriteIndexAcquire(uint64_t value) = 0;
/// @brief Updates the Write index using Relaxed semantics
///
/// @param value Value of new write index
///
/// @return uint64_t Value of write index before the update
virtual uint64_t AddWriteIndexRelaxed(uint64_t value) = 0;
/// @brief Updates the Write index using Release semantics
///
/// @param value Value of new write index
///
/// @return uint64_t Value of write index before the update
virtual uint64_t AddWriteIndexRelease(uint64_t value) = 0;
/// @brief Set CU Masking
///
/// @param num_cu_mask_count size of mask bit array
///
/// @param cu_mask pointer to cu mask
///
/// @return hsa_status_t
virtual hsa_status_t SetCUMasking(const uint32_t num_cu_mask_count,
const uint32_t* cu_mask) = 0;
// Handle of AMD Queue struct
amd_queue_t& amd_queue_;
hsa_queue_t* public_handle() const { return public_handle_; }
protected:
static void set_public_handle(Queue* ptr, hsa_queue_t* handle) {
ptr->do_set_public_handle(handle);
}
virtual void do_set_public_handle(hsa_queue_t* handle) {
public_handle_ = handle;
}
hsa_queue_t* public_handle_;
private:
DISALLOW_COPY_AND_ASSIGN(Queue);
};
}
#endif // header guard
+204
View File
@@ -0,0 +1,204 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
// This file is used only for open source cmake builds, if we hardcode the
// register values in amd_aql_queue.cpp then this file won't be required. For
// now we are using this file where register details are spelled out in the
// structs/unions below.
#ifndef HSA_RUNTME_CORE_INC_REGISTERS_H_
#define HSA_RUNTME_CORE_INC_REGISTERS_H_
typedef enum SQ_RSRC_BUF_TYPE {
SQ_RSRC_BUF = 0x00000000,
SQ_RSRC_BUF_RSVD_1 = 0x00000001,
SQ_RSRC_BUF_RSVD_2 = 0x00000002,
SQ_RSRC_BUF_RSVD_3 = 0x00000003,
} SQ_RSRC_BUF_TYPE;
typedef enum BUF_DATA_FORMAT {
BUF_DATA_FORMAT_INVALID = 0x00000000,
BUF_DATA_FORMAT_8 = 0x00000001,
BUF_DATA_FORMAT_16 = 0x00000002,
BUF_DATA_FORMAT_8_8 = 0x00000003,
BUF_DATA_FORMAT_32 = 0x00000004,
BUF_DATA_FORMAT_16_16 = 0x00000005,
BUF_DATA_FORMAT_10_11_11 = 0x00000006,
BUF_DATA_FORMAT_11_11_10 = 0x00000007,
BUF_DATA_FORMAT_10_10_10_2 = 0x00000008,
BUF_DATA_FORMAT_2_10_10_10 = 0x00000009,
BUF_DATA_FORMAT_8_8_8_8 = 0x0000000a,
BUF_DATA_FORMAT_32_32 = 0x0000000b,
BUF_DATA_FORMAT_16_16_16_16 = 0x0000000c,
BUF_DATA_FORMAT_32_32_32 = 0x0000000d,
BUF_DATA_FORMAT_32_32_32_32 = 0x0000000e,
BUF_DATA_FORMAT_RESERVED_15 = 0x0000000f,
} BUF_DATA_FORMAT;
typedef enum BUF_NUM_FORMAT {
BUF_NUM_FORMAT_UNORM = 0x00000000,
BUF_NUM_FORMAT_SNORM = 0x00000001,
BUF_NUM_FORMAT_USCALED = 0x00000002,
BUF_NUM_FORMAT_SSCALED = 0x00000003,
BUF_NUM_FORMAT_UINT = 0x00000004,
BUF_NUM_FORMAT_SINT = 0x00000005,
BUF_NUM_FORMAT_SNORM_OGL__SI__CI = 0x00000006,
BUF_NUM_FORMAT_RESERVED_6__VI = 0x00000006,
BUF_NUM_FORMAT_FLOAT = 0x00000007,
} BUF_NUM_FORMAT;
typedef enum SQ_SEL_XYZW01 {
SQ_SEL_0 = 0x00000000,
SQ_SEL_1 = 0x00000001,
SQ_SEL_RESERVED_0 = 0x00000002,
SQ_SEL_RESERVED_1 = 0x00000003,
SQ_SEL_X = 0x00000004,
SQ_SEL_Y = 0x00000005,
SQ_SEL_Z = 0x00000006,
SQ_SEL_W = 0x00000007,
} SQ_SEL_XYZW01;
union COMPUTE_TMPRING_SIZE {
struct {
#if defined(LITTLEENDIAN_CPU)
unsigned int WAVES : 12;
unsigned int WAVESIZE : 13;
unsigned int : 7;
#elif defined(BIGENDIAN_CPU)
unsigned int : 7;
unsigned int WAVESIZE : 13;
unsigned int WAVES : 12;
#endif
} bitfields, bits;
unsigned int u32All;
signed int i32All;
float f32All;
};
union SQ_BUF_RSRC_WORD0 {
struct {
#if defined(LITTLEENDIAN_CPU)
unsigned int BASE_ADDRESS : 32;
#elif defined(BIGENDIAN_CPU)
unsigned int BASE_ADDRESS : 32;
#endif
} bitfields, bits;
unsigned int u32All;
signed int i32All;
float f32All;
};
union SQ_BUF_RSRC_WORD1 {
struct {
#if defined(LITTLEENDIAN_CPU)
unsigned int BASE_ADDRESS_HI : 16;
unsigned int STRIDE : 14;
unsigned int CACHE_SWIZZLE : 1;
unsigned int SWIZZLE_ENABLE : 1;
#elif defined(BIGENDIAN_CPU)
unsigned int SWIZZLE_ENABLE : 1;
unsigned int CACHE_SWIZZLE : 1;
unsigned int STRIDE : 14;
unsigned int BASE_ADDRESS_HI : 16;
#endif
} bitfields, bits;
unsigned int u32All;
signed int i32All;
float f32All;
};
union SQ_BUF_RSRC_WORD2 {
struct {
#if defined(LITTLEENDIAN_CPU)
unsigned int NUM_RECORDS : 32;
#elif defined(BIGENDIAN_CPU)
unsigned int NUM_RECORDS : 32;
#endif
} bitfields, bits;
unsigned int u32All;
signed int i32All;
float f32All;
};
union SQ_BUF_RSRC_WORD3 {
struct {
#if defined(LITTLEENDIAN_CPU)
unsigned int DST_SEL_X : 3;
unsigned int DST_SEL_Y : 3;
unsigned int DST_SEL_Z : 3;
unsigned int DST_SEL_W : 3;
unsigned int NUM_FORMAT : 3;
unsigned int DATA_FORMAT : 4;
unsigned int ELEMENT_SIZE : 2;
unsigned int INDEX_STRIDE : 2;
unsigned int ADD_TID_ENABLE : 1;
unsigned int ATC__CI__VI : 1;
unsigned int HASH_ENABLE : 1;
unsigned int HEAP : 1;
unsigned int MTYPE__CI__VI : 3;
unsigned int TYPE : 2;
#elif defined(BIGENDIAN_CPU)
unsigned int TYPE : 2;
unsigned int MTYPE__CI__VI : 3;
unsigned int HEAP : 1;
unsigned int HASH_ENABLE : 1;
unsigned int ATC__CI__VI : 1;
unsigned int ADD_TID_ENABLE : 1;
unsigned int INDEX_STRIDE : 2;
unsigned int ELEMENT_SIZE : 2;
unsigned int DATA_FORMAT : 4;
unsigned int NUM_FORMAT : 3;
unsigned int DST_SEL_W : 3;
unsigned int DST_SEL_Z : 3;
unsigned int DST_SEL_Y : 3;
unsigned int DST_SEL_X : 3;
#endif
} bitfields, bits;
unsigned int u32All;
signed int i32All;
float f32All;
};
#endif // header guard
+498
View File
@@ -0,0 +1,498 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
// HSA runtime C++ interface file.
#ifndef HSA_RUNTME_CORE_INC_RUNTIME_H_
#define HSA_RUNTME_CORE_INC_RUNTIME_H_
#include <vector>
#include <map>
#include "core/inc/hsa_ext_interface.h"
#include "core/inc/hsa_internal.h"
#include "core/inc/agent.h"
#include "core/inc/memory_region.h"
#include "core/inc/signal.h"
#include "core/util/utils.h"
#include "core/util/locks.h"
#include "core/util/os.h"
#include "core/inc/amd_loader_context.hpp"
#include "amd_hsa_code.hpp"
//---------------------------------------------------------------------------//
// Constants //
//---------------------------------------------------------------------------//
#define HSA_ARGUMENT_ALIGN_BYTES 16
#define HSA_QUEUE_ALIGN_BYTES 64
#define HSA_PACKET_ALIGN_BYTES 64
namespace core {
extern bool g_use_interrupt_wait;
/// @brief Runtime class provides the following functions:
/// - open and close connection to kernel driver.
/// - load supported extension library (image and finalizer).
/// - load tools library.
/// - expose supported agents.
/// - allocate and free memory.
/// - memory copy and fill.
/// - grant access to memory (dgpu memory pool extension).
/// - maintain loader state.
/// - monitor asynchronous event from agent.
class Runtime {
public:
/// @brief Structure to describe connectivity between agents.
struct LinkInfo {
uint32_t num_hop;
hsa_amd_memory_pool_link_info_t info;
};
/// @brief Open connection to kernel driver and increment reference count.
/// @retval True if the connection to kernel driver is successfully opened.
static bool Acquire();
/// @brief Checks if connection to kernel driver is opened.
/// @retval True if the connection to kernel driver is opened.
static bool IsOpen();
// @brief Callback handler for VM fault access.
static bool VMFaultHandler(hsa_signal_value_t val, void* arg);
/// @brief Singleton object of the runtime.
static Runtime* runtime_singleton_;
/// @brief Decrement reference count and close connection to kernel driver.
/// @retval True if reference count is larger than 0.
bool Release();
/// @brief Insert agent into agent list ::agents_.
/// @param [in] agent Pointer to the agent object.
void RegisterAgent(Agent* agent);
/// @brief Delete all agent objects from ::agents_.
void DestroyAgents();
/// @brief Set the number of links connecting the agents in the platform.
void SetLinkCount(size_t num_link);
/// @brief Register link information connecting @p node_id_from and @p
/// node_id_to.
/// @param [in] node_id_from Node id of the source node.
/// @param [in] node_id_to Node id of the destination node.
/// @param [in] link_info The link information between source and destination
/// nodes.
void RegisterLinkInfo(uint32_t node_id_from, uint32_t node_id_to,
uint32_t num_hop,
hsa_amd_memory_pool_link_info_t& link_info);
/// @brief Query link information between two nodes.
/// @param [in] node_id_from Node id of the source node.
/// @param [in] node_id_to Node id of the destination node.
/// @retval The link information between source and destination nodes.
const LinkInfo& GetLinkInfo(uint32_t node_id_from, uint32_t node_id_to);
/// @brief Invoke the user provided call back for each agent in the agent
/// list.
///
/// @param [in] callback User provided callback function.
/// @param [in] data User provided pointer as input for @p callback.
///
/// @retval ::HSA_STATUS_SUCCESS if the callback function for each traversed
/// agent returns ::HSA_STATUS_SUCCESS.
hsa_status_t IterateAgent(hsa_status_t (*callback)(hsa_agent_t agent,
void* data),
void* data);
/// @brief Allocate memory on a particular region.
///
/// @param [in] region Pointer to region object.
/// @param [in] size Allocation size in bytes.
/// @param [out] address Pointer to store the allocation result.
///
/// @retval ::HSA_STATUS_SUCCESS If allocation is successful.
hsa_status_t AllocateMemory(const MemoryRegion* region, size_t size,
void** address);
/// @brief Allocate memory on a particular region with option to restrict
/// access to the owning agent.
///
/// @param [in] restrict_access If true, the allocation result would only be
/// accessible to the agent(s) that own the region object.
/// @param [in] region Pointer to region object.
/// @param [in] size Allocation size in bytes.
/// @param [out] address Pointer to store the allocation result.
///
/// @retval ::HSA_STATUS_SUCCESS If allocation is successful.
hsa_status_t AllocateMemory(bool restrict_access, const MemoryRegion* region,
size_t size, void** address);
/// @brief Free memory previously allocated with AllocateMemory.
///
/// @param [in] ptr Address of the memory to be freed.
///
/// @retval ::HSA_STATUS_ERROR If @p ptr is not the address of previous
/// allocation via ::core::Runtime::AllocateMemory
/// @retval ::HSA_STATUS_SUCCESS if @p ptr is successfully released.
hsa_status_t FreeMemory(void* ptr);
/// @brief Blocking memory copy from src to dst.
///
/// @param [in] dst Memory address of the destination.
/// @param [in] src Memory address of the source.
/// @param [in] size Copy size in bytes.
///
/// @retval ::HSA_STATUS_SUCCESS if memory copy is successful and completed.
hsa_status_t CopyMemory(void* dst, const void* src, size_t size);
/// @brief Non-blocking memory copy from src to dst.
///
/// @details The memory copy will be performed after all signals in
/// @p dep_signals have value of 0. On completion @p completion_signal
/// will be decremented.
///
/// @param [in] dst Memory address of the destination.
/// @param [in] dst_agent Agent object associated with the destination. This
/// agent should be able to access the destination and source.
/// @param [in] src Memory address of the source.
/// @param [in] src_agent Agent object associated with the source. This
/// agent should be able to access the destination and source.
/// @param [in] size Copy size in bytes.
/// @param [in] dep_signals Array of signal dependency.
/// @param [in] completion_signal Completion signal object.
///
/// @retval ::HSA_STATUS_SUCCESS if copy command has been submitted
/// successfully to the agent DMA queue.
hsa_status_t CopyMemory(void* dst, core::Agent& dst_agent, const void* src,
core::Agent& src_agent, size_t size,
std::vector<core::Signal*>& dep_signals,
core::Signal& completion_signal);
/// @brief Fill the first @p count of uint32_t in ptr with value.
///
/// @param [in] ptr Memory address to be filled.
/// @param [in] value The value/pattern that will be used to set @p ptr.
/// @param [in] count Number of uint32_t element to be set.
///
/// @retval ::HSA_STATUS_SUCCESS if memory fill is successful and completed.
hsa_status_t FillMemory(void* ptr, uint32_t value, size_t count);
/// @brief Set agents as the whitelist to access ptr.
///
/// @param [in] num_agents The number of agent handles in @p agents array.
/// @param [in] agents Agent handle array.
/// @param [in] ptr Pointer of memory previously allocated via
/// core::Runtime::AllocateMemory.
///
/// @retval ::HSA_STATUS_SUCCESS The whitelist has been configured
/// successfully and all agents in the @p agents could start accessing @p ptr.
hsa_status_t AllowAccess(uint32_t num_agents, const hsa_agent_t* agents,
const void* ptr);
/// @brief Query system information.
///
/// @param [in] attribute System info attribute to query.
/// @param [out] value Pointer to store the attribute value.
///
/// @retval HSA_STATUS_SUCCESS The attribute is valid and the @p value is
/// set.
hsa_status_t GetSystemInfo(hsa_system_info_t attribute, void* value);
/// @brief Query next available queue id.
///
/// @retval Next available queue id.
uint32_t GetQueueId();
/// @brief Register a callback function @p handler that is associated with
/// @p signal to asynchronous event monitor thread.
///
/// @param [in] signal Signal handle associated with @p handler.
/// @param [in] cond The condition to execute the @p handler.
/// @param [in] value The value to compare with @p signal value. If the
/// comparison satisfy @p cond, the @p handler will be called.
/// @param [in] arg Pointer to the argument that will be provided to @p
/// handler.
///
/// @retval ::HSA_STATUS_SUCCESS Registration is successful.
hsa_status_t SetAsyncSignalHandler(hsa_signal_t signal,
hsa_signal_condition_t cond,
hsa_signal_value_t value,
hsa_amd_signal_handler handler, void* arg);
hsa_status_t InteropMap(uint32_t num_agents, Agent** agents,
int interop_handle, uint32_t flags, size_t* size,
void** ptr, size_t* metadata_size,
const void** metadata);
hsa_status_t InteropUnmap(void* ptr);
const std::vector<Agent*>& cpu_agents() { return cpu_agents_; }
const std::vector<Agent*>& gpu_agents() { return gpu_agents_; }
const std::vector<uint32_t>& gpu_ids() { return gpu_ids_; }
Agent* blit_agent() { return blit_agent_; }
Agent* host_agent() { return host_agent_; }
const std::vector<const MemoryRegion*>& system_regions_fine() const {
return system_regions_fine_;
}
const std::vector<const MemoryRegion*>& system_regions_coarse() const {
return system_regions_coarse_;
}
amd::hsa::loader::Loader* loader() { return loader_; }
amd::LoaderContext* loader_context() { return &loader_context_; }
amd::hsa::code::AmdHsaCodeManager* code_manager() { return &code_manager_; }
std::function<void*(size_t, size_t)>& system_allocator() {
return system_allocator_;
}
std::function<void(void*)>& system_deallocator() {
return system_deallocator_;
}
ExtensionEntryPoints extensions_;
protected:
static void AsyncEventsLoop(void*);
struct AllocationRegion {
AllocationRegion() : region(NULL), assigned_agent_(NULL), size(0) {}
AllocationRegion(const MemoryRegion* region_arg, size_t size_arg)
: region(region_arg), assigned_agent_(NULL), size(size_arg) {}
const MemoryRegion* region;
const Agent* assigned_agent_;
size_t size;
};
struct AsyncEventsControl {
AsyncEventsControl() : async_events_thread_(NULL) {}
void Shutdown();
hsa_signal_t wake;
os::Thread async_events_thread_;
KernelMutex lock;
bool exit;
};
struct AsyncEvents {
void PushBack(hsa_signal_t signal, hsa_signal_condition_t cond,
hsa_signal_value_t value, hsa_amd_signal_handler handler,
void* arg);
void CopyIndex(size_t dst, size_t src);
size_t Size();
void PopBack();
void Clear();
std::vector<hsa_signal_t> signal_;
std::vector<hsa_signal_condition_t> cond_;
std::vector<hsa_signal_value_t> value_;
std::vector<hsa_amd_signal_handler> handler_;
std::vector<void*> arg_;
};
// Will be created before any user could call hsa_init but also could be
// destroyed before incorrectly written programs call hsa_shutdown.
static KernelMutex bootstrap_lock_;
Runtime();
Runtime(const Runtime&);
Runtime& operator=(const Runtime&);
~Runtime() {}
/// @brief Open connection to kernel driver.
void Load();
/// @brief Close connection to kernel driver and cleanup resources.
void Unload();
/// @brief Dynamically load extension libraries (images, finalizer) and
/// call OnLoad method on each loaded library.
void LoadExtensions();
/// @brief Call OnUnload method on each extension library then close it.
void UnloadExtensions();
/// @brief Dynamically load tool libraries and call OnUnload method on each
/// loaded library.
void LoadTools();
/// @brief Call OnUnload method of each tool library.
void UnloadTools();
/// @brief Close tool libraries.
void CloseTools();
// @brief Binds virtual memory access fault handler to this node.
void BindVmFaultHandler();
/// @brief Blocking memory copy from src to dst. One of the src or dst
/// is user pointer. A particular setup need to be made if the DMA queue
/// for the memory copy belongs to a dGPU agent. E.g: pin the user pointer
/// before copying, or using a staging buffer.
///
/// @param [in] dst Memory address of the destination.
/// @param [in] src Memory address of the source.
/// @param [in] size Copy size in bytes.
/// @param [in] dst_malloc If true, then @p dst is the user pointer. Otherwise
/// @p src is the user pointer.
///
/// @retval ::HSA_STATUS_SUCCESS if memory copy is successful and completed.
hsa_status_t CopyMemoryHostAlloc(void* dst, const void* src, size_t size,
bool dst_malloc);
/// @brief Get the index of ::link_matrix_.
/// @param [in] node_id_from Node id of the source node.
/// @param [in] node_id_to Node id of the destination node.
/// @retval Index in ::link_matrix_.
uint32_t GetIndexLinkInfo(uint32_t node_id_from, uint32_t node_id_to);
// Mutex object to protect multithreaded access to ::Acquire and ::Release.
KernelMutex kernel_lock_;
// Mutex object to protect multithreaded access to ::allocation_map_.
KernelMutex memory_lock_;
// Array containing tools library handles.
std::vector<os::LibHandle> tool_libs_;
// Agent list containing all CPU agents in the platform.
std::vector<Agent*> cpu_agents_;
// Agent list containing all compatible GPU agents in the platform.
std::vector<Agent*> gpu_agents_;
// Agent list containing all compatible gpu agent ids in the platform.
std::vector<uint32_t> gpu_ids_;
// List of all fine grain system memory region in the platform.
std::vector<const MemoryRegion*> system_regions_fine_;
// List of all coarse grain system memory region in the platform.
std::vector<const MemoryRegion*> system_regions_coarse_;
// Matrix of IO link.
std::vector<LinkInfo> link_matrix_;
// Loader instance.
amd::hsa::loader::Loader* loader_;
// Loader context.
amd::LoaderContext loader_context_;
// Code object manager.
amd::hsa::code::AmdHsaCodeManager code_manager_;
// Contains the region, address, and size of previously allocated memory.
std::map<const void*, AllocationRegion> allocation_map_;
// Allocator using ::system_region_
std::function<void*(size_t, size_t)> system_allocator_;
// Deallocator using ::system_region_
std::function<void(void*)> system_deallocator_;
// Pointer to a host/cpu agent object.
Agent* host_agent_;
// Pointer to DMA agent.
Agent* blit_agent_;
AsyncEventsControl async_events_control_;
AsyncEvents async_events_;
AsyncEvents new_async_events_;
// Queue id counter.
uint32_t queue_count_;
// Starting address of SVM address space.
// On APU the cpu and gpu could access the area inside starting and end of
// the SVM address space.
// On dGPU, only the gpu is guaranteed to have access to the area inside the
// SVM address space, since it maybe backed by private gpu VRAM.
uintptr_t start_svm_address_;
// End address of SVM address space.
// start_svm_address_ + size
uintptr_t end_svm_address_;
// System clock frequency.
uint64_t sys_clock_freq_;
// @brief AMD HSA event to monitor for virtual memory access fault.
HsaEvent* vm_fault_event_;
// @brief HSA signal to contain the VM fault event.
Signal* vm_fault_signal_;
// Holds reference count to runtime object.
volatile uint32_t ref_count_;
// Frees runtime memory when the runtime library is unloaded if safe to do so.
// Failure to release the runtime indicates an incorrect application but is
// common (example: calls library routines at process exit).
friend class RuntimeCleanup;
};
} // namespace core
#endif // header guard
+269
View File
@@ -0,0 +1,269 @@
////////////////////////////////////////////////////////////////////////////////
//
// The University of Illinois/NCSA
// Open Source License (NCSA)
//
// Copyright (c) 2014-2015, Advanced Micro Devices, Inc. All rights reserved.
//
// Developed by:
//
// AMD Research and AMD HSA Software Development
//
// Advanced Micro Devices, Inc.
//
// www.amd.com
//
// Permission is hereby granted, free of charge, to any person obtaining a copy
// of this software and associated documentation files (the "Software"), to
// deal with 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:
//
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimers.
// - Redistributions in binary form must reproduce the above copyright
// notice, this list of conditions and the following disclaimers in
// the documentation and/or other materials provided with the distribution.
// - Neither the names of Advanced Micro Devices, Inc,
// nor the names of its contributors may be used to endorse or promote
// products derived from this Software without specific prior written
// permission.
//
// 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 CONTRIBUTORS 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 WITH THE SOFTWARE.
//
////////////////////////////////////////////////////////////////////////////////
// HSA runtime C++ interface file.
#ifndef HSA_RUNTME_CORE_INC_SIGNAL_H_
#define HSA_RUNTME_CORE_INC_SIGNAL_H_
#include "hsakmt.h"
#include "core/common/shared.h"
#include "core/inc/runtime.h"
#include "core/inc/checked.h"
#include "core/util/utils.h"
#include "inc/amd_hsa_signal.h"
namespace core {
class Signal;
/// @brief Helper structure to simplify conversion of amd_signal_t and
/// core::Signal object.
struct SharedSignal {
amd_signal_t amd_signal;
Signal* core_signal;
};
/// @brief An abstract base class which helps implement the public hsa_signal_t
/// type (an opaque handle) and its associated APIs. At its core, signal uses
/// a 32 or 64 bit value. This value can be waitied on or signaled atomically
/// using specified memory ordering semantics.
class Signal : public Checked<0x71FCCA6A3D5D5276>,
public Shared<SharedSignal, AMD_SIGNAL_ALIGN_BYTES> {
public:
/// @brief Constructor initializes the signal with initial value.
explicit Signal(hsa_signal_value_t initial_value)
: Shared(), signal_(shared_object()->amd_signal) {
if (!Shared::IsSharedObjectAllocationValid()) {
invalid_ = true;
return;
}
shared_object()->core_signal = this;
signal_.kind = AMD_SIGNAL_KIND_INVALID;
signal_.value = initial_value;
invalid_ = false;
waiting_ = 0;
retained_ = 0;
}
virtual ~Signal() { signal_.kind = AMD_SIGNAL_KIND_INVALID; }
bool IsValid() const {
if (CheckedType::IsValid() && !invalid_) return true;
return false;
}
/// @brief Converts from this implementation class to the public
/// hsa_signal_t type - an opaque handle.
static __forceinline hsa_signal_t Convert(Signal* signal) {
const uint64_t handle =
(signal != NULL && signal->IsValid())
? static_cast<uint64_t>(
reinterpret_cast<uintptr_t>(&signal->signal_))
: 0;
const hsa_signal_t signal_handle = {handle};
return signal_handle;
}
/// @brief Converts from this implementation class to the public
/// hsa_signal_t type - an opaque handle.
static __forceinline const hsa_signal_t Convert(const Signal* signal) {
const uint64_t handle =
(signal != NULL && signal->IsValid())
? static_cast<uint64_t>(
reinterpret_cast<uintptr_t>(&signal->signal_))
: 0;
const hsa_signal_t signal_handle = {handle};
return signal_handle;
}
/// @brief Converts from public hsa_signal_t type (an opaque handle) to
/// this implementation class object.
static __forceinline Signal* Convert(hsa_signal_t signal) {
return (signal.handle != 0)
? reinterpret_cast<const SharedSignal*>(
static_cast<uintptr_t>(signal.handle) -
(reinterpret_cast<uintptr_t>(
&reinterpret_cast<SharedSignal*>(1234)->amd_signal) -
uintptr_t(1234)))->core_signal
: NULL;
}
// Below are various methods corresponding to the APIs, which load/store the
// signal value or modify the existing signal value automically and with
// specified memory ordering semantics.
virtual hsa_signal_value_t LoadRelaxed() = 0;
virtual hsa_signal_value_t LoadAcquire() = 0;
virtual void StoreRelaxed(hsa_signal_value_t value) = 0;
virtual void StoreRelease(hsa_signal_value_t value) = 0;
virtual hsa_signal_value_t WaitRelaxed(hsa_signal_condition_t condition,
hsa_signal_value_t compare_value,
uint64_t timeout,
hsa_wait_state_t wait_hint) = 0;
virtual hsa_signal_value_t WaitAcquire(hsa_signal_condition_t condition,
hsa_signal_value_t compare_value,
uint64_t timeout,
hsa_wait_state_t wait_hint) = 0;
virtual void AndRelaxed(hsa_signal_value_t value) = 0;
virtual void AndAcquire(hsa_signal_value_t value) = 0;
virtual void AndRelease(hsa_signal_value_t value) = 0;
virtual void AndAcqRel(hsa_signal_value_t value) = 0;
virtual void OrRelaxed(hsa_signal_value_t value) = 0;
virtual void OrAcquire(hsa_signal_value_t value) = 0;
virtual void OrRelease(hsa_signal_value_t value) = 0;
virtual void OrAcqRel(hsa_signal_value_t value) = 0;
virtual void XorRelaxed(hsa_signal_value_t value) = 0;
virtual void XorAcquire(hsa_signal_value_t value) = 0;
virtual void XorRelease(hsa_signal_value_t value) = 0;
virtual void XorAcqRel(hsa_signal_value_t value) = 0;
virtual void AddRelaxed(hsa_signal_value_t value) = 0;
virtual void AddAcquire(hsa_signal_value_t value) = 0;
virtual void AddRelease(hsa_signal_value_t value) = 0;
virtual void AddAcqRel(hsa_signal_value_t value) = 0;
virtual void SubRelaxed(hsa_signal_value_t value) = 0;
virtual void SubAcquire(hsa_signal_value_t value) = 0;
virtual void SubRelease(hsa_signal_value_t value) = 0;
virtual void SubAcqRel(hsa_signal_value_t value) = 0;
virtual hsa_signal_value_t ExchRelaxed(hsa_signal_value_t value) = 0;
virtual hsa_signal_value_t ExchAcquire(hsa_signal_value_t value) = 0;
virtual hsa_signal_value_t ExchRelease(hsa_signal_value_t value) = 0;
virtual hsa_signal_value_t ExchAcqRel(hsa_signal_value_t value) = 0;
virtual hsa_signal_value_t CasRelaxed(hsa_signal_value_t expected,
hsa_signal_value_t value) = 0;
virtual hsa_signal_value_t CasAcquire(hsa_signal_value_t expected,
hsa_signal_value_t value) = 0;
virtual hsa_signal_value_t CasRelease(hsa_signal_value_t expected,
hsa_signal_value_t value) = 0;
virtual hsa_signal_value_t CasAcqRel(hsa_signal_value_t expected,
hsa_signal_value_t value) = 0;
//-------------------------
// implementation specific
//-------------------------
typedef void* rtti_t;
/// @brief Returns the address of the value.
virtual hsa_signal_value_t* ValueLocation() const = 0;
/// @brief Applies only to InterrupEvent type, returns the event used to.
/// Returns NULL for DefaultEvent Type.
virtual HsaEvent* EopEvent() = 0;
/// @brief Waits until any signal in the list satisfies its condition or
/// timeout is reached.
/// Returns the index of a satisfied signal. Returns -1 on timeout and
/// errors.
static uint32_t WaitAny(uint32_t signal_count, hsa_signal_t* hsa_signals,
hsa_signal_condition_t* conds,
hsa_signal_value_t* values, uint64_t timeout_hint,
hsa_wait_state_t wait_hint,
hsa_signal_value_t* satisfying_value);
__forceinline bool IsType(rtti_t id) { return _IsA(id); }
/// @brief Allows special case interaction with signal destruction cleanup.
void Retain() { atomic::Increment(&retained_); }
void Release() { atomic::Decrement(&retained_); }
/// @brief Checks if signal is currently in use such that it should not be
/// deleted.
bool InUse() const { return (retained_ != 0) || (waiting_ != 0); }
/// @brief Checks if signal is currently in use by a wait API.
bool InWaiting() const { return waiting_ != 0; }
/// @brief Structure which defines key signal elements like type and value.
/// Address of this struct is used as a value for the opaque handle of type
/// hsa_signal_t provided to the public API.
amd_signal_t& signal_;
protected:
/// @brief Simple RTTI type checking helper
/// Returns true if the object can be converted to the query type via
/// static_cast.
/// Do not use directly. Use IsType in the desired derived type instead.
virtual bool _IsA(rtti_t id) const = 0;
/// @variable Indicates if signal is valid or not.
volatile bool invalid_;
/// @variable Indicates number of runtime threads waiting on this signal.
/// Value of zero means no waits.
volatile uint32_t waiting_;
volatile uint32_t retained_;
private:
DISALLOW_COPY_AND_ASSIGN(Signal);
};
struct hsa_signal_handle {
hsa_signal_t signal;
hsa_signal_handle() {}
hsa_signal_handle(hsa_signal_t Signal) { signal = Signal; }
operator hsa_signal_t() { return signal; }
Signal* operator->() { return core::Signal::Convert(signal); }
};
static_assert(
sizeof(hsa_signal_handle) == sizeof(hsa_signal_t),
"hsa_signal_handle and hsa_signal_t must have identical binary layout.");
static_assert(
sizeof(hsa_signal_handle[2]) == sizeof(hsa_signal_t[2]),
"hsa_signal_handle and hsa_signal_t must have identical binary layout.");
} // namespace core
#endif // header guard