blob: 5a9bbd588cbd0db589309deaadf268e0d5608a68 [file] [log] [blame]
/*
* Copyright (c) 2018 Advanced Micro Devices, Inc.
* All rights reserved.
*
* For use for simulation and test purposes only
*
* Redistribution and use in source and binary forms, with or without
* modification, are permitted provided that the following conditions are met:
*
* 1. Redistributions of source code must retain the above copyright notice,
* this list of conditions and the following disclaimer.
*
* 2. Redistributions in binary form must reproduce the above copyright notice,
* this list of conditions and the following disclaimer in the documentation
* and/or other materials provided with the distribution.
*
* 3. Neither the name of the copyright holder nor the names of its
* contributors may be used to endorse or promote products derived from this
* software without specific prior written permission.
*
* THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
* AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
* IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
* ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
* LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
* CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
* SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
* INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
* CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
* ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
* POSSIBILITY OF SUCH DAMAGE.
*/
#include "gpu-compute/gpu_command_processor.hh"
#include "debug/GPUCommandProc.hh"
#include "debug/GPUKernelInfo.hh"
#include "gpu-compute/dispatcher.hh"
#include "params/GPUCommandProcessor.hh"
#include "sim/process.hh"
#include "sim/proxy_ptr.hh"
#include "sim/syscall_emul_buf.hh"
GPUCommandProcessor::GPUCommandProcessor(const Params &p)
: HSADevice(p), dispatcher(*p.dispatcher), driver(nullptr)
{
dispatcher.setCommandProcessor(this);
}
/**
* submitDispatchPkt() is the entry point into the CP from the HSAPP
* and is only meant to be used with AQL kernel dispatch packets.
* After the HSAPP receives and extracts an AQL packet, it sends
* it to the CP, which is responsible for gathering all relevant
* information about a task, initializing CU state, and sending
* it to the dispatcher for WG creation and dispatch.
*
* First we need capture all information from the the AQL pkt and
* the code object, then store it in an HSAQueueEntry. Once the
* packet and code are extracted, we extract information from the
* queue descriptor that the CP needs to perform state initialization
* on the CU. Finally we call dispatch() to send the task to the
* dispatcher. When the task completely finishes, we call finishPkt()
* on the HSA packet processor in order to remove the packet from the
* queue, and notify the runtime that the task has completed.
*/
void
GPUCommandProcessor::submitDispatchPkt(void *raw_pkt, uint32_t queue_id,
Addr host_pkt_addr)
{
static int dynamic_task_id = 0;
_hsa_dispatch_packet_t *disp_pkt = (_hsa_dispatch_packet_t*)raw_pkt;
/**
* we need to read a pointer in the application's address
* space to pull out the kernel code descriptor.
*/
auto *tc = sys->threads[0];
auto &virt_proxy = tc->getVirtProxy();
/**
* The kernel_object is a pointer to the machine code, whose entry
* point is an 'amd_kernel_code_t' type, which is included in the
* kernel binary, and describes various aspects of the kernel. The
* desired entry is the 'kernel_code_entry_byte_offset' field,
* which provides the byte offset (positive or negative) from the
* address of the amd_kernel_code_t to the start of the machine
* instructions.
*/
AMDKernelCode akc;
virt_proxy.readBlob(disp_pkt->kernel_object, (uint8_t*)&akc,
sizeof(AMDKernelCode));
DPRINTF(GPUCommandProc, "GPU machine code is %lli bytes from start of the "
"kernel object\n", akc.kernel_code_entry_byte_offset);
DPRINTF(GPUCommandProc,"GPUCommandProc: Sending dispatch pkt to %lu\n",
(uint64_t)tc->cpuId());
Addr machine_code_addr = (Addr)disp_pkt->kernel_object
+ akc.kernel_code_entry_byte_offset;
DPRINTF(GPUCommandProc, "Machine code starts at addr: %#x\n",
machine_code_addr);
Addr kern_name_addr(0);
std::string kernel_name;
/**
* BLIT kernels don't have symbol names. BLIT kernels are built-in compute
* kernels issued by ROCm to handle DMAs for dGPUs when the SDMA
* hardware engines are unavailable or explicitly disabled. They can also
* be used to do copies that ROCm things would be better performed
* by the shader than the SDMA engines. They are also sometimes used on
* APUs to implement asynchronous memcopy operations from 2 pointers in
* host memory. I have no idea what BLIT stands for.
* */
if (akc.runtime_loader_kernel_symbol) {
virt_proxy.readBlob(akc.runtime_loader_kernel_symbol + 0x10,
(uint8_t*)&kern_name_addr, 0x8);
virt_proxy.readString(kernel_name, kern_name_addr);
} else {
kernel_name = "Blit kernel";
}
DPRINTF(GPUKernelInfo, "Kernel name: %s\n", kernel_name.c_str());
HSAQueueEntry *task = new HSAQueueEntry(kernel_name, queue_id,
dynamic_task_id, raw_pkt, &akc, host_pkt_addr, machine_code_addr);
DPRINTF(GPUCommandProc, "Task ID: %i Got AQL: wg size (%dx%dx%d), "
"grid size (%dx%dx%d) kernarg addr: %#x, completion "
"signal addr:%#x\n", dynamic_task_id, disp_pkt->workgroup_size_x,
disp_pkt->workgroup_size_y, disp_pkt->workgroup_size_z,
disp_pkt->grid_size_x, disp_pkt->grid_size_y,
disp_pkt->grid_size_z, disp_pkt->kernarg_address,
disp_pkt->completion_signal);
DPRINTF(GPUCommandProc, "Extracted code object: %s (num vector regs: %d, "
"num scalar regs: %d, code addr: %#x, kernarg size: %d, "
"LDS size: %d)\n", kernel_name, task->numVectorRegs(),
task->numScalarRegs(), task->codeAddr(), 0, 0);
initABI(task);
++dynamic_task_id;
}
uint64_t
GPUCommandProcessor::functionalReadHsaSignal(Addr signal_handle)
{
Addr value_addr = getHsaSignalValueAddr(signal_handle);
auto tc = system()->threads[0];
ConstVPtr<Addr> prev_value(value_addr, tc);
return *prev_value;
}
void
GPUCommandProcessor::updateHsaSignal(Addr signal_handle, uint64_t signal_value)
{
// The signal value is aligned 8 bytes from
// the actual handle in the runtime
Addr value_addr = getHsaSignalValueAddr(signal_handle);
Addr mailbox_addr = getHsaSignalMailboxAddr(signal_handle);
Addr event_addr = getHsaSignalEventAddr(signal_handle);
DPRINTF(GPUCommandProc, "Triggering completion signal: %x!\n", value_addr);
Addr *new_signal = new Addr;
*new_signal = signal_value;
dmaWriteVirt(value_addr, sizeof(Addr), nullptr, new_signal, 0);
auto tc = system()->threads[0];
ConstVPtr<uint64_t> mailbox_ptr(mailbox_addr, tc);
// Notifying an event with its mailbox pointer is
// not supported in the current implementation. Just use
// mailbox pointer to distinguish between interruptible
// and default signal. Interruptible signal will have
// a valid mailbox pointer.
if (*mailbox_ptr != 0) {
// This is an interruptible signal. Now, read the
// event ID and directly communicate with the driver
// about that event notification.
ConstVPtr<uint32_t> event_val(event_addr, tc);
DPRINTF(GPUCommandProc, "Calling signal wakeup event on "
"signal event value %d\n", *event_val);
signalWakeupEvent(*event_val);
}
}
void
GPUCommandProcessor::attachDriver(HSADriver *hsa_driver)
{
fatal_if(driver, "Should not overwrite driver.");
driver = hsa_driver;
}
/**
* submitVendorPkt() is for accepting vendor-specific packets from
* the HSAPP. Vendor-specific packets may be used by the runtime to
* send commands to the HSA device that are specific to a particular
* vendor. The vendor-specific packets should be defined by the vendor
* in the runtime.
*/
/**
* TODO: For now we simply tell the HSAPP to finish the packet,
* however a future patch will update this method to provide
* the proper handling of any required vendor-specific packets.
* In the version of ROCm that is currently supported (1.6)
* the runtime will send packets that direct the CP to
* invalidate the GPUs caches. We do this automatically on
* each kernel launch in the CU, so this is safe for now.
*/
void
GPUCommandProcessor::submitVendorPkt(void *raw_pkt, uint32_t queue_id,
Addr host_pkt_addr)
{
hsaPP->finishPkt(raw_pkt, queue_id);
}
/**
* submitAgentDispatchPkt() is for accepting agent dispatch packets.
* These packets will control the dispatch of Wg on the device, and inform
* the host when a specified number of Wg have been executed on the device.
*
* For now it simply finishes the pkt.
*/
void
GPUCommandProcessor::submitAgentDispatchPkt(void *raw_pkt, uint32_t queue_id,
Addr host_pkt_addr)
{
//Parse the Packet, see what it wants us to do
_hsa_agent_dispatch_packet_t * agent_pkt =
(_hsa_agent_dispatch_packet_t *)raw_pkt;
if (agent_pkt->type == AgentCmd::Nop) {
DPRINTF(GPUCommandProc, "Agent Dispatch Packet NOP\n");
} else if (agent_pkt->type == AgentCmd::Steal) {
//This is where we steal the HSA Task's completion signal
int kid = agent_pkt->arg[0];
DPRINTF(GPUCommandProc,
"Agent Dispatch Packet Stealing signal handle for kernel %d\n",
kid);
HSAQueueEntry *task = dispatcher.hsaTask(kid);
uint64_t signal_addr = task->completionSignal();// + sizeof(uint64_t);
uint64_t return_address = agent_pkt->return_address;
DPRINTF(GPUCommandProc, "Return Addr: %p\n",return_address);
//*return_address = signal_addr;
Addr *new_signal_addr = new Addr;
*new_signal_addr = (Addr)signal_addr;
dmaWriteVirt(return_address, sizeof(Addr), nullptr, new_signal_addr, 0);
DPRINTF(GPUCommandProc,
"Agent Dispatch Packet Stealing signal handle from kid %d :" \
"(%x:%x) writing into %x\n",
kid,signal_addr,new_signal_addr,return_address);
} else
{
panic("The agent dispatch packet provided an unknown argument in" \
"arg[0],currently only 0(nop) or 1(return kernel signal) is accepted");
}
hsaPP->finishPkt(raw_pkt, queue_id);
}
/**
* Once the CP has finished extracting all relevant information about
* a task and has initialized the ABI state, we send a description of
* the task to the dispatcher. The dispatcher will create and dispatch
* WGs to the CUs.
*/
void
GPUCommandProcessor::dispatchPkt(HSAQueueEntry *task)
{
dispatcher.dispatch(task);
}
void
GPUCommandProcessor::signalWakeupEvent(uint32_t event_id)
{
driver->signalWakeupEvent(event_id);
}
/**
* The CP is responsible for traversing all HSA-ABI-related data
* structures from memory and initializing the ABI state.
* Information provided by the MQD, AQL packet, and code object
* metadata will be used to initialze register file state.
*/
void
GPUCommandProcessor::initABI(HSAQueueEntry *task)
{
auto *readDispIdOffEvent = new ReadDispIdOffsetDmaEvent(*this, task);
Addr hostReadIdxPtr
= hsaPP->getQueueDesc(task->queueId())->hostReadIndexPtr;
dmaReadVirt(hostReadIdxPtr + sizeof(hostReadIdxPtr),
sizeof(readDispIdOffEvent->readDispIdOffset), readDispIdOffEvent,
&readDispIdOffEvent->readDispIdOffset);
}
System*
GPUCommandProcessor::system()
{
return sys;
}
AddrRangeList
GPUCommandProcessor::getAddrRanges() const
{
AddrRangeList ranges;
return ranges;
}
void
GPUCommandProcessor::setShader(Shader *shader)
{
_shader = shader;
}
Shader*
GPUCommandProcessor::shader()
{
return _shader;
}