ROCM Kernel生成到GPU执行全流程分析
·
ROCm Kernel执行流程深度分析
目录
概述
本文档详细分析ROCm软件栈中从MIOpen到GPU硬件的完整kernel执行流程，包括kernel编译、加载、入队、AQL包生成、队列管理和Graph API机制。
软件栈层级
┌─────────────────────────────────────────────────────────────┐
│ Layer 1: MIOpen (深度学习库) │
│ - Kernel Cache管理 │
│ - Solver选择和算法实现 │
├─────────────────────────────────────────────────────────────┤
│ Layer 2: HIP Runtime (CLR) │
│ - Module加载 │
│ - Kernel启动API │
├─────────────────────────────────────────────────────────────┤
│ Layer 3: ROCclr (虚拟设备层) │
│ - Command Queue管理 │
│ - AQL包生成和分发 │
├─────────────────────────────────────────────────────────────┤
│ Layer 4: ROCR Runtime (HSA Runtime) │
│ - Queue管理 │
│ - Signal/Doorbell机制 │
├─────────────────────────────────────────────────────────────┤
│ Layer 5: ROCT-Thunk-Interface │
│ - KFD驱动接口 │
├─────────────────────────────────────────────────────────────┤
│ Layer 6: KFD Kernel Driver │
│ - GPU硬件管理 │
└─────────────────────────────────────────────────────────────┘
Kernel生成与编译
完整编译调用链
miopenPoolingForward (API Entry)
│
v
PoolingDescriptor::Forward
│
v
SolverContainer::ExecutePrimitive
│
v
SearchForSolutions → FindSolution → Solver::GetSolution
│
v
Handle::PrepareInvoker → KernelCache::AddKernel
│
v
Handle::LoadProgram → HIPOCProgram → COMgr/HIPRTC Compilation
│
v
Binary Cache Storage/Retrieval
1.1 API入口: miopenPoolingForward
文件: MIOpen/src/pooling_api.cpp:312-340
extern "C" miopenStatus_t miopenPoolingForward(miopenHandle_t handle,
const miopenPoolingDescriptor_t poolDesc,
const void* alpha,
const miopenTensorDescriptor_t xDesc,
const void* x,
const void* beta,
const miopenTensorDescriptor_t yDesc,
void* y,
bool do_backward,
void* workSpace,
size_t workSpaceSize)
{
return miopen::try_([&] {
miopen::deref(poolDesc).Forward(miopen::deref(handle),
alpha,
miopen::deref(xDesc),
DataCast(x),
beta,
miopen::deref(yDesc),
DataCast(y),
do_backward,
DataCast(workSpace),
workSpaceSize);
});
}
1.2 PoolingDescriptor::Forward实现
文件: MIOpen/src/ocl/pooling_ocl.cpp:57-140
miopenStatus_t PoolingDescriptor::Forward(Handle& handle,
const void* alpha,
const TensorDescriptor& xDesc,
ConstData_t x,
const void* beta,
const TensorDescriptor& yDesc,
Data_t y,
bool save_index,
Data_t workSpace,
size_t workSpaceSize) const
{
// 创建问题描述
const auto algo_name = AlgorithmName{"miopenPoolingForwardDirect"};
const auto problem = pooling::ProblemDescription{*this, xDesc, yDesc, save_index};
// 创建调用参数
const auto invoke_params = [&]() {
auto tmp = pooling::FwdInvokeParams{};
tmp.type = InvokeType::Run;
tmp.xDesc = xDesc;
tmp.yDesc = yDesc;
tmp.pooling = *this;
tmp.x = x;
tmp.y = y;
tmp.workspace = workSpace;
tmp.workspace_size = workSpaceSize;
return tmp;
}();
// 执行solver
PoolingForwardSolvers().ExecutePrimitive(handle, problem, algo_name, invoke_params);
}
1.3 Solver容器定义
文件: MIOpen/src/ocl/pooling_ocl.cpp:40-47
static auto PoolingForwardSolvers()
{
return solver::SolverContainer<
solver::pooling::PoolingForward2d, // 优化的2D pooling
solver::pooling::PoolingForwardNd, // 通用N-D pooling
solver::pooling::PoolingForwardNaive, // 简单动态solver
solver::pooling::TransposedPoolingFwd2d,// Layout转换wrapper
solver::pooling::TransposedPoolingFwdNd // Layout转换wrapper
>{};
}
1.4 ExecutePrimitive流程
文件: MIOpen/src/include/miopen/find_solution.hpp:476-503
template <class Problem>
void ExecutePrimitive(const ExecutionContext& ctx,
const Problem& problem,
const AlgorithmName& algo,
const AnyInvokeParams& invoke_params) const
{
const auto network_config = problem.MakeNetworkConfig();
// 1. 检查是否已有缓存的invoker
if(const auto existingInvoker =
ctx.GetStream().GetInvoker(network_config, std::nullopt, algo))
{
(*existingInvoker)(ctx.GetStream(), invoke_params);
return;
}
// 2. 搜索适用的solution (限制为1个)
const auto slns = SearchForSolutions(ctx, problem, 1, invoke_params);
if(slns.empty())
MIOPEN_THROW(miopenStatusNotImplemented, "No solver found.");
// 3. 获取第一个solution并准备invoker
const auto& sln = slns.front();
const auto invoker =
ctx.GetStream().PrepareInvoker(*sln.invoker_factory, sln.construction_params);
// 4. 注册invoker供后续使用并执行
ctx.GetStream().RegisterInvoker(invoker, network_config, sln.solver_id, algo);
invoker(ctx.GetStream(), invoke_params);
}
1.5 Solver选择: SearchForSolutions
文件: MIOpen/src/include/miopen/find_solution.hpp:331-395
template <class Problem, class Solution = miopen::solver::ConvSolution>
std::vector<Solution>
SearchForSolutions(const ExecutionContext& ctx,
const Problem& problem,
std::size_t limit = std::numeric_limits<std::size_t>::max(),
const AnyInvokeParams& invoke_params = {}) const
{
std::vector<Solution> ss;
std::size_t count = 0;
miopen::each_args(
[&](auto solver) {
if(count >= limit)
return;
// 检查solver是否适用
if(!solver.IsApplicable(ctx, problem))
{
MIOPEN_LOG_I2(solver.SolverDbId() << ": Not applicable");
}
else
{
// 找到适用的solver，获取solution
auto s = FindSolution(solver, ctx, problem, db, invoke_params, "", std::nullopt);
if(s.Succeeded())
{
++count;
ss.emplace_back(std::move(s));
}
}
},
Solvers{}...); // 展开所有solver
return ss;
}
1.6 Solver实现示例: PoolingForward2d
文件: MIOpen/src/solver/pooling/forward2d.cpp:136-262
IsApplicable检查
bool PoolingForward2d::IsApplicable(const ExecutionContext& context,
const miopen::pooling::ProblemDescription& problem) const
{
return problem.GetDirection() == miopen::pooling::Direction::Forward &&
problem.GetXDesc().GetNumDims() == 4 && // 4D tensor (NCHW)
problem.GetXDesc().GetType() == problem.GetYDesc().GetType() &&
(problem.GetXDesc().GetType() == miopenFloat ||
problem.GetXDesc().GetType() == miopenHalf) &&
problem.GetXDesc().IsPossibleLayout4D5D("NCHW") &&
problem.GetYDesc().IsPossibleLayout4D5D("NCHW") &&
sizeof_private_memory(problem) <=
TargetProperties::GetMaxWaveScratchSize() / context.GetStream().GetWavefrontWidth();
}
GetSolution - Kernel构建
ConvSolution PoolingForward2d::GetSolution(const ExecutionContext&,
const miopen::pooling::ProblemDescription& problem) const
{
auto result = ConvSolution{miopenStatusSuccess};
{
auto kernel = KernelInfo{};
kernel.kernel_file = "MIOpenPooling.cl"; // 源文件
kernel.kernel_name = "mloPoolingG"; // 入口函数
// 计算kernel参数
const kernel_params kp(problem);
// 构建编译参数
auto build_params = KernelBuildParameters{
{"MLO_POOLING_OP_ID", pooling_method},
{"MLO_POOLING_KERNEL_SZ1", kp.kernel_size_h},
{"MLO_POOLING_STRIDE1", kp.kernel_stride_h},
{"MLO_POOLING_KERNEL_SZ0", kp.kernel_size_w},
{"MLO_POOLING_STRIDE0", kp.kernel_stride_w},
{"MLO_POOLING_N_HORIZ_OUT_PIX", kp.out_pix_tile0},
{"MLO_POOLING_N_VERT_OUT_PIX", kp.out_pix_tile1},
{"MLO_POOLING_GROUP_SZ0", grp_tile0},
{"MLO_POOLING_GROUP_SZ1", grp_tile1},
{"MLO_POOLING_INDEX_TYPE", get_pooling_index_type_name(pool_d.GetIndexType())},
};
kernel.comp_options = build_params.GenerateFor(kbp::OpenCL{});
// 设置workgroup维度
kernel.l_wk.push_back(grp_tile0);
kernel.l_wk.push_back(grp_tile1);
kernel.l_wk.push_back(1);
// 设置global work维度
kernel.g_wk.push_back(static_cast<std::size_t>(g_wk_width) * grp_tile0);
kernel.g_wk.push_back(static_cast<std::size_t>(g_wk_height) * grp_tile1);
kernel.g_wk.push_back(static_cast<std::size_t>(n_outputs) * batch_sz);
result.construction_params.push_back(kernel);
}
// 定义invoker factory - 创建实际的kernel调用
result.invoker_factory = [](const std::vector<Kernel>& kernels) {
return [=](const Handle& handle_, const AnyInvokeParams& raw_params) {
decltype(auto) kernel = handle_.Run(kernels.front());
decltype(auto) params = raw_params.CastTo<miopen::pooling::FwdInvokeParams>();
kernel(params.x, params.y, params.workspace, /* ... args ... */);
};
};
return result;
}
1.7 Handle::PrepareInvoker和KernelCache
文件: MIOpen/src/hip/handlehip.cpp:484-514
Invoker Handle::PrepareInvoker(const InvokerFactory& factory,
const std::vector<solver::KernelInfo>& kernels,
std::vector<Program>* programs_out) const
{
std::vector<Kernel> built;
built.reserve(kernels.size());
if(programs_out != nullptr)
programs_out->resize(kernels.size());
for(auto i = 0; i < kernels.size(); ++i)
{
const auto& k = kernels[i];
Program* program_out = programs_out != nullptr ? &(*programs_out)[i] : nullptr;
MIOPEN_LOG_I2("Preparing kernel: " << k.kernel_name);
// 添加kernel到cache
const auto kernel = this->impl->cache.AddKernel(*this,
"",
"",
k.kernel_file,
k.kernel_name,
k.l_wk,
k.g_wk,
k.comp_options,
kernels.size(),
"",
program_out);
built.push_back(kernel);
}
return factory(built);
}
1.8 KernelCache::AddKernel - Program加载
文件: MIOpen/src/kernel_cache.cpp:94-155
Kernel KernelCache::AddKernel(const Handle& h,
const std::string& algorithm,
const std::string& network_config,
const fs::path& program_name,
const std::string& kernel_name,
const std::vector<size_t>& vld,
const std::vector<size_t>& vgd,
std::string params,
std::size_t cache_index,
const std::string& kernel_src,
Program* program_out)
{
const auto program = [&] {
// 检查program是否已在cache中
auto program_it = program_map.find(std::make_pair(program_name, params));
if(program_it != program_map.end())
{
auto& program = program_it->second;
return program;
}
else
{
// 加载/编译program
auto program = h.LoadProgram(program_name, params, kernel_src, program_out != nullptr);
program_map[std::make_pair(program_name, params)] = program;
return program;
}
}();
// 从program创建kernel
Kernel kernel{program, kernel_name, vld, vgd};
// 如果提供了network config，缓存kernel
if(!network_config.empty() && !algorithm.empty())
{
this->AddKernel(key, kernel, cache_index);
}
return kernel;
}
1.9 Handle::LoadProgram - 核心编译流程
文件: MIOpen/src/hip/handlehip.cpp:536-652
Program Handle::LoadProgram(const fs::path& program_name,
std::string params,
const std::string& kernel_src,
bool force_attach_binary) const
{
this->impl->set_ctx();
std::string arch_name = this->GetTargetProperties().Name();
std::string orig_params = params;
// 添加目标架构到编译参数
params += " -mcpu=" + this->GetTargetProperties().Name();
// 步骤1: 尝试从binary cache加载
auto hsaco = miopen::LoadBinary(
this->GetTargetProperties(), this->GetMaxComputeUnits(), program_name, params);
// 如果目标ID特定的binary未找到，回退到基础架构
if(hsaco.empty())
{
const auto arch_target_id = miopen::SplitDelim(arch_name, ':');
if(arch_target_id.size() > 1)
{
const auto base_arch = arch_target_id.at(0);
hsaco = miopen::LoadBinary(this->GetTargetProperties(),
this->GetMaxComputeUnits(),
program_name,
orig_params + " -mcpu=" + base_arch);
}
}
// 步骤2: 如果不在cache中，从源码编译
if(hsaco.empty())
{
CompileTimer ct;
auto p = HIPOCProgram{program_name.string(), params,
this->GetTargetProperties(), kernel_src};
ct.Log("Kernel", program_name.string());
// 步骤3: 保存到cache
#if MIOPEN_ENABLE_SQLITE_KERN_CACHE
std::vector<char> binary;
if(!p.IsCodeObjectInMemory())
binary = miopen::LoadFile(p.GetCodeObjectPathname());
miopen::SaveBinary(p.IsCodeObjectInMemory() ? p.GetCodeObjectBlob() : binary,
this->GetTargetProperties(),
this->GetMaxComputeUnits(),
program_name,
params);
#endif
return p;
}
else
{
// 使用cached binary
auto p = HIPOCProgram{program_name, hsaco};
return p;
}
}
1.10 HIPOCProgram和COMgr编译
文件: MIOpen/src/hipoc/hipoc_program.cpp:193-329
HIPOCProgram构造函数
HIPOCProgramImpl::HIPOCProgramImpl(const fs::path& program_name,
std::string params,
const TargetProperties& target_,
const std::string& kernel_src)
: program(program_name), target(target_)
{
BuildCodeObject(params, kernel_src);
if(!binary.empty())
{
module = CreateModuleInMem(binary);
}
else
{
module = CreateModule(hsaco_file);
}
}
BuildCodeObject with COMgr
void HIPOCProgramImpl::BuildCodeObject(std::string params, const std::string& kernel_src)
{
const auto src = [&]() -> std::string_view {
if(program.extension() == ".mlir")
return {};
if(!kernel_src.empty())
return kernel_src;
return GetKernelSrc(program); // 获取embedded source
}();
#if MIOPEN_BUILD_DEV
if(program.extension() == ".cpp")
params += " -Werror" + HipKernelWarningsString();
else if(program.extension() == ".cl")
params += " -Werror" + OclKernelWarningsString();
#else
if(program.extension() == ".cpp" || program.extension() == ".cl")
params += " -Wno-everything";
#endif
#if MIOPEN_USE_COMGR
BuildCodeObjectInMemory(params, src, program); // 使用COMgr
#else
BuildCodeObjectInFile(params, src, program); // 使用offline compiler
#endif
}
COMgr内存构建
void HIPOCProgramImpl::BuildCodeObjectInMemory(const std::string& params,
const std::string_view src,
const fs::path& filename)
{
if(filename.extension() == dynamic_library_postfix) // ".so"或".dll"
{
binary.resize(src.size());
std::memcpy(&binary[0], src.data(), src.size());
}
else
{
#if MIOPEN_WORKAROUND_ROCM_COMPILER_SUPPORT_ISSUE_27
static std::mutex mutex;
std::lock_guard<std::mutex> lock(mutex);
#endif
if(filename.extension() == ".cpp")
{
hiprtc::BuildHip(filename.string(), src, params, target, binary);
}
else if(filename.extension() == ".s")
{
comgr::BuildAsm(filename.string(), src, params, target, binary);
}
else
{
comgr::BuildOcl(filename.string(), src, params, target, binary);
}
}
}
1.11 COMgr集成细节
文件: MIOpen/src/comgr.cpp:597-648
OpenCL Build via COMgr (BuildOcl)
void BuildOcl(const std::string& name,
std::string_view text,
const std::string& options,
const miopen::TargetProperties& target,
std::vector<char>& binary)
{
PrintVersion();
try
{
const Dataset inputs;
inputs.AddData(name, text, AMD_COMGR_DATA_KIND_SOURCE);
const ActionInfo action;
action.SetLanguage(AMD_COMGR_LANGUAGE_OPENCL_2_0);
SetIsaName(action, target, true);
action.SetLogging(true);
// 准备编译器选项
auto optCompile = miopen::SplitSpaceSeparated(options);
compiler::lc::ocl::RemoveOptionsUnwanted(optCompile);
compiler::lc::ocl::AddCompilerOptions(optCompile);
action.SetOptionList(optCompile);
// COMgr action pipeline:
const Dataset addedPch;
action.Do(AMD_COMGR_ACTION_ADD_PRECOMPILED_HEADERS, inputs, addedPch);
const Dataset linkedBc;
action.Do(AMD_COMGR_ACTION_COMPILE_SOURCE_WITH_DEVICE_LIBS_TO_BC, addedPch, linkedBc);
action.SetOptionList(optCompile);
const Dataset relocatable;
action.Do(AMD_COMGR_ACTION_CODEGEN_BC_TO_RELOCATABLE, linkedBc, relocatable);
action.SetOptionList(OptionList());
const Dataset exe;
action.Do(AMD_COMGR_ACTION_LINK_RELOCATABLE_TO_EXECUTABLE, relocatable, exe);
// 提取binary
const auto data = exe.GetData(AMD_COMGR_DATA_KIND_EXECUTABLE, 0);
data.GetBytes(binary);
}
catch(ComgrError& ex)
{
binary.resize(0);
MIOPEN_LOG_E("comgr status = " << GetStatusText(ex));
}
}
1.12 Binary Cache系统
文件: MIOpen/src/binary_cache.cpp:152-196
SQLite-based Cache
#if MIOPEN_ENABLE_SQLITE_KERN_CACHE
std::vector<char> LoadBinary(const TargetProperties& target,
const size_t num_cu,
const fs::path& name,
const std::string& args)
{
if(miopen::IsCacheDisabled())
return {};
auto db = GetDb(target, num_cu);
const auto filename = make_object_file_name(name);
const KernelConfig cfg{filename, args, {}};
MIOPEN_LOG_I2("Loading binary for: " << filename << "; args: " << args);
auto record = db.FindRecord(cfg);
if(record)
{
MIOPEN_LOG_I2("Successfully loaded binary for: " << filename);
return *record;
}
else
{
MIOPEN_LOG_I2("Unable to load binary for: " << filename);
return {};
}
}
void SaveBinary(const std::vector<char>& hsaco,
const TargetProperties& target,
const std::size_t num_cu,
const fs::path& name,
const std::string& args)
{
if(miopen::IsCacheDisabled())
return;
auto db = GetDb(target, num_cu);
const auto filename = make_object_file_name(name);
KernelConfig cfg{filename, args, hsaco};
MIOPEN_LOG_I2("Saving binary for: " << filename << "; args: " << args);
db.StoreRecord(cfg);
}
#endif
1.13 Kernel源文件嵌入
文件: MIOpen/src/kernels/kernel.cpp.in
Kernel源码在build时通过code generation template嵌入到binary中：
// 生成kernel source lookup的template
std::string_view GetKernelSrc(const fs::path& name)
{
static const std::unordered_map<fs::path, std::string_view, FsPathHash> data{
${INIT_KERNELS} // build时生成
};
auto it = kernels().find(name.filename());
if(it == kernels().end())
MIOPEN_THROW("Failed to load kernel source: " + name.filename());
return it->second;
}
Pooling相关Kernel文件:
MIOpen/src/kernels/MIOpenPooling.cl- 2D pooling kernelMIOpen/src/kernels/MIOpenPoolingND.cl- ND pooling kernelMIOpen/src/kernels/MIOpenPoolingForwardNaive.cl- Naive实现
1.14 MIOpen Kernel Cache系统
文件: MIOpen/src/kernel_cache.cpp
MIOpen使用两级缓存系统：
class KernelCache {
using Key = std::pair<std::string, std::string>; // (algorithm, network_config)
using KernelMap = std::unordered_map<Key, std::vector<Kernel>, SimpleHash>;
using ProgramMap = std::unordered_map<Key, Program, SimpleHash>;
KernelMap kernel_map_; // Kernel缓存
ProgramMap program_map_; // Program缓存
};
缓存层级:
- Invoker Cache: 缓存已准备好的invoker，避免重复solver选择
- Kernel Cache: 缓存已编译的kernel对象
- Program Cache: 缓存已加载的HIP program
- Binary Cache: 持久化的SQLite数据库，存储编译后的binary
**编译步骤**:
1. 使用COMgr (Code Object Manager) 编译kernel源代码
2. 生成HIP二进制代码对象
3. 调用 `hipModuleLoadData` 加载module
### 1.3 Module加载
**文件**: `MIOpen/src/hipoc/hipoc_program.cpp:166-181`
```cpp
template <typename T>
hipModulePtr CreateModuleInMem(const T& blob) {
hipModule_t raw_m;
// 加载module到HIP runtime
auto status = hipModuleLoadData(&raw_m, reinterpret_cast<const void*>(blob.data()));
hipModulePtr m{raw_m};
return m;
}
HIP Runtime中的Module加载:
文件: clr/hipamd/src/hip_module.cpp
hipError_t hipModuleLoadData(hipModule_t* module, const void* image) {
// 解析代码对象
// 提取kernel符号
// 创建module对象
}
Kernel入队机制
2.1 Command Queue结构
文件: clr/rocclr/platform/commandqueue.cpp
class HostQueue : public CommandQueue {
public:
HostQueue(Context& context, Device& device, cl_command_queue_properties props,
uint queueRTCUs, Priority priority, const std::vector<uint32_t>& cuMask)
: CommandQueue(context, device, props, device.info().queueProperties_,
queueRTCUs, priority, cuMask),
lastEnqueueCommand_(nullptr),
head_(nullptr),
tail_(nullptr),
isActive_(false) {
if (AMD_DIRECT_DISPATCH) {
thread_.Init(this); // 直接分发模式
} else {
thread_.start(this); // 线程模式
}
}
private:
Command* head_; // 队列头
Command* tail_; // 队列尾
Command* lastEnqueueCommand_;
};
2.2 Command入队流程
文件: clr/rocclr/platform/command.cpp
void Command::enqueue() {
if (AMD_DIRECT_DISPATCH) {
// 直接分发模式
setStatus(CL_QUEUED);
// 通知等待事件
for (const auto& event : eventWaitList()) {
event->notifyCmdQueue(!kCpuWait);
}
// 提交batch
ScopedLock sl(queue_->vdev()->execution());
queue_->FormSubmissionBatch(this);
// 提交到设备
submit(*queue_->vdev());
queue_->FlushSubmissionBatch(this);
} else {
// 线程模式
queue_->append(*this);
queue_->flush();
}
}
2.3 Kernel命令创建
文件: clr/hipamd/src/hip_module.cpp:357-436
hipError_t ihipModuleLaunchKernel(hipFunction_t f,
uint32_t globalWorkSizeX, ...,
hipStream_t stream,
void** kernelParams,
void** extra, ...) {
// 获取kernel对象
amd::Kernel* kernel = reinterpret_cast<amd::Kernel*>(f);
// 设置kernel参数
if (kernelParams != nullptr) {
for (size_t i = 0; i < kernel->signature().numParameters(); ++i) {
kernel->parameters().set(i, kernelParams[i]);
}
}
// 创建NDRangeKernelCommand
amd::NDRangeKernelCommand* kernelCommand = new amd::NDRangeKernelCommand(
*stream, // 目标stream
waitList, // 等待事件列表
*kernel, // kernel对象
ndrange, // 执行范围
sharedMemBytes, // 共享内存大小
params, // kernel参数
...
);
// 入队命令
kernelCommand->enqueue();
}
AQL包详解
3.1 AQL Packet结构
文件: ROCR-Runtime/runtime/hsa-runtime/inc/hsa.h
AQL (Architected Queueing Language) 是AMD GPU的硬件队列协议。
Kernel Dispatch Packet (64字节)
typedef struct hsa_kernel_dispatch_packet_s {
uint16_t header; // 包头 (类型、barrier、fence scope)
uint16_t setup; // 设置 (维度)
uint16_t workgroup_size_x; // workgroup X维度
uint16_t workgroup_size_y; // workgroup Y维度
uint16_t workgroup_size_z; // workgroup Z维度
uint16_t reserved0;
uint32_t grid_size_x; // grid X维度
uint32_t grid_size_y; // grid Y维度
uint32_t grid_size_z; // grid Z维度
uint32_t private_segment_size; // 私有内存大小
uint32_t group_segment_size; // LDS大小
uint64_t kernel_object; // kernel代码对象句柄
void* kernarg_address; // kernel参数地址
uint64_t reserved2;
hsa_signal_t completion_signal; // 完成信号
} hsa_kernel_dispatch_packet_t;
Packet Header格式 (16位)
Bit 0-7: Packet Type
- 1: Kernel Dispatch
- 2: Barrier-AND
- 3: Barrier-OR
- 4: Agent Dispatch
Bit 8: Barrier (1=后续包需等待此包完成)
Bit 9-10: Acquire Fence Scope
- 0: None
- 1: Agent
- 2: System
Bit 11-12: Release Fence Scope
- 0: None
- 1: Agent
- 2: System
3.2 AQL包生成流程
文件: clr/rocclr/device/rocm/rocvirtual.cpp:2993-3389
bool VirtualGPU::submitKernelInternal(...) {
// 1. 获取GPU Kernel对象
const ROCKernel* devKernel = static_cast<const ROCKernel*>(&kernel);
const ROCDevice& dev = static_cast<const ROCDevice&>(kernel.program().device());
// 2. 获取kernel代码句柄
const GpuKernel& gpuKernel = devKernel->gpuKernel(dev);
// 3. 准备kernel参数
const auto& params = kernel.parameters();
size_t ldsUsage = kernel.workGroupInfo()->localMemSize_;
// 4. 分配参数缓冲区
void* argBuffer = allocKernArg(devKernel->kernargSegmentByteSize(),
devKernel->kernargSegmentAlignment());
// 5. 填充参数
char* writeAddress = static_cast<char*>(argBuffer);
for (size_t i = 0; i < params.size(); ++i) {
if (params[i].size_ != 0) {
memcpy(writeAddress, params[i].data_, params[i].size_);
writeAddress += params[i].size_;
}
}
// 6. 创建AQL Packet
hsa_kernel_dispatch_packet_t dispatchPacket{};
// 初始化为无效header，防止GPU提前处理
dispatchPacket.header = kInvalidAql;
// 设置kernel代码对象
dispatchPacket.kernel_object = gpuKernel.KernelCodeHandle();
// 设置执行维度
dispatchPacket.grid_size_x = global[0];
dispatchPacket.grid_size_y = global[1];
dispatchPacket.grid_size_z = global[2];
dispatchPacket.workgroup_size_x = local[0];
dispatchPacket.workgroup_size_y = local[1];
dispatchPacket.workgroup_size_z = local[2];
// 设置参数地址和内存大小
dispatchPacket.kernarg_address = argBuffer;
dispatchPacket.group_segment_size = ldsUsage + sharedMemBytes;
dispatchPacket.private_segment_size = devKernel->workGroupInfo()->privateMemSize_;
// 7. 分发AQL包
if (!dispatchAqlPacket(&dispatchPacket, aqlHeaderWithOrder,
(sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS),
GPU_FLUSH_ON_EXECUTION, false, nullptr, attach_signal)) {
return false;
}
return true;
}
3.3 AQL包分发
文件: clr/rocclr/device/rocm/rocvirtual.cpp:840-951
template <typename AqlPacket>
bool VirtualGPU::dispatchGenericAqlPacket(
AqlPacket* packet, uint16_t header, uint16_t rest,
bool blocking, bool attach_signal) {
const uint32_t queueSize = gpu_queue_->size;
const uint32_t queueMask = queueSize - 1;
const uint32_t sw_queue_size = queueMask;
// 1. 原子获取写索引
uint64_t index = hsa_queue_add_write_index_screlease(gpu_queue_, 1);
// 2. 处理fence scope
if (addSystemScope_) {
header &= ~(HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE |
HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
header |= (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE |
HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
addSystemScope_ = false;
}
// 3. 获取完成信号
packet->completion_signal = Barriers().ActiveSignal(kInitSignalValueOne,
timestamp_, attachSignal);
// 4. 等待队列槽位可用
while ((index - hsa_queue_load_read_index_scacquire(gpu_queue_)) >= sw_queue_size) {
amd::Os::yield();
}
// 5. 写入AQL包到队列内存
TrackQueueProgress(*packet, index);
AqlPacket* aql_loc = &((AqlPacket*)(gpu_queue_->base_address))[index & queueMask];
*aql_loc = *packet;
// 6. 写入header (带release语义)
if (header != 0) {
packet_store_release(reinterpret_cast<uint32_t*>(aql_loc), header, rest);
}
// 7. 触发Doorbell通知GPU
hsa_signal_store_screlease(gpu_queue_->doorbell_signal, index);
// 8. 标记有未完成的dispatch
hasPendingDispatch_ = true;
// 9. 如果需要阻塞等待
if (blocking) {
if (!Barriers().WaitCurrent()) {
return false;
}
}
return true;
}
队列中的命令流动
4.1 Queue内存布局
GPU Queue内存结构 (环形缓冲区)
┌─────────────────────────────────────────────────────────────────┐
│ Queue Base Address (gpu_queue_->base_address) │
├─────────────────────────────────────────────────────────────────┤
│ Index 0 │ Index 1 │ Index 2 │ ... │ Index N-1 │
│ (64 bytes) │ (64 bytes) │ (64 bytes) │ │ (64 bytes) │
├─────────────────────────────────────────────────────────────────┤
│ AQL Packet │ AQL Packet │ AQL Packet │ │ AQL Packet │
└─────────────────────────────────────────────────────────────────┘
读写索引管理:
- Read Index: GPU CP (Command Processor) 读取位置
- Write Index: CPU写入位置
- Queue Size: 必须是2的幂 (如 256, 512, 1024)
4.2 命令状态流转
Command状态机:
CL_QUEUED ──────▶ CL_SUBMITTED ──────▶ CL_RUNNING ──────▶ CL_COMPLETE
│ │ │ │
│ │ │ │
入队完成 提交到GPU GPU开始执行 执行完成
(enqueue) (submit) (CP读取) (signal)
4.3 批量提交优化
文件: clr/rocclr/device/rocm/rocvirtual.cpp
void VirtualGPU::FormSubmissionBatch(Command* command) {
// 将多个命令批量提交，减少doorbell次数
if (batch_head_ == nullptr) {
batch_head_ = command;
}
batch_tail_ = command;
batch_size_++;
}
void VirtualGPU::FlushSubmissionBatch(Command* last_command) {
if (batch_size_ > 0) {
// 一次性提交所有命令
// 只触发一次doorbell
hsa_signal_store_screlease(gpu_queue_->doorbell_signal, write_index);
batch_size_ = 0;
batch_head_ = nullptr;
batch_tail_ = nullptr;
}
}
Doorbell机制
5.1 Doorbell原理
Doorbell是CPU通知GPU有新工作的机制，基于内存映射的信号量。
Doorbell机制流程:
CPU侧: GPU侧 (Command Processor):
┌─────────────┐ ┌─────────────────┐
│ 1. 写入AQL │ │ 3. 检测Doorbell │
│ Packet │ │ 变化 │
│ 到队列 │ │ │
└──────┬──────┘ └────────┬────────┘
│ │
▼ ▼
┌─────────────┐ ┌─────────────────┐
│ 2. 写入 │ ──────Doorbell Signal────▶│ 4. 从队列读取 │
│ Doorbell │ (内存映射) │ AQL Packet │
│ (MMIO) │ │ │
└─────────────┘ └─────────────────┘
│
▼
┌─────────────────┐
│ 5. 解析并执行 │
│ Kernel │
└─────────────────┘
5.2 Doorbell类型
文件: ROCR-Runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp:461-535
void AqlQueue::StoreRelaxed(hsa_signal_value_t value) {
if (doorbell_type_ == 2) {
// Type 2: 硬件Doorbell，支持AQL语义
atomic::Store(signal_.hardware_doorbell_ptr, uint64_t(value),
std::memory_order_release);
} else {
// Type 1: 传统Doorbell
uint32_t legacy_dispatch_id = uint32_t(value);
if (legacy_dispatch_id != 0) {
// 转换为硬件队列ID
legacy_dispatch_id = (legacy_dispatch_id << 2) + 1;
}
atomic::Store(signal_.legacy_hardware_doorbell_ptr, legacy_dispatch_id,
std::memory_order_release);
}
}
5.3 Queue创建和Doorbell映射
文件: ROCR-Runtime/libhsakmt/src/queues.c:700-751
HSAKMT_STATUS HSAKMTAPI hsaKmtCreateQueue(
HSAuint32 NodeId,
HSA_QUEUE_TYPE Type,
HSAuint32 QueuePercentage,
HSA_QUEUE_PRIORITY Priority,
void* QueueAddress,
HSAuint64 QueueSizeInBytes,
HsaEvent* Event,
HsaQueueResource* QueueResource) {
// 设置队列资源
HsaQueueResource queue_rsrc = {0};
queue_rsrc.Queue_read_ptr_aql = (uint64_t*)&amd_queue_.read_dispatch_id;
queue_rsrc.Queue_write_ptr_aql = (uint64_t*)&amd_queue_.write_dispatch_id;
// 调用KFD创建队列
err = hsakmt_ioctl(hsakmt_kfd_fd, AMDKFD_IOC_CREATE_QUEUE, &args);
// 映射Doorbell内存
err = map_doorbell(NodeId, gpu_id, doorbell_mmap_offset);
// 返回Doorbell指针给调用者
QueueResource->Queue_DoorBell = VOID_PTR_ADD(
doorbells[NodeId].mapping, doorbell_offset);
return HSAKMT_STATUS_SUCCESS;
}
Graph API原理
6.1 Graph API概述
Graph API允许捕获一系列操作并在后续重放，减少CPU开销。
Graph API优势:
- 减少CPU开销: 预先生成AQL包，避免运行时重复构建
- 优化依赖: 静态分析依赖关系，优化执行顺序
- 支持整图优化: 可以跨kernel进行优化
6.2 Graph节点结构
文件: clr/hipamd/src/hip_graph_internal.hpp
class GraphNode : public hipGraphNodeDOTAttribute {
public:
GraphNode(hipGraphNodeType type, const char* style = "",
const char* shape = "", const char* label = "")
: type_(type),
visited_(false),
inDegree_(0),
outDegree_(0),
id_(nextID++),
parentGraph_(nullptr),
isEnabled_(1) {}
// 捕获并形成AQL包
hipError_t CaptureAndFormPacket(GraphKernelArgManager* kernArgMgr) {
// 获取捕获stream
auto capture_stream = hip::getNullStream(
g_devices[dev_id_]->devices()[0]->context(), false);
// 创建command
hipError_t status = CreateCommand(capture_stream);
// 清除之前的包
gpuPackets_.clear();
// 捕获GPU Packet
for (auto& command : commands_) {
command->setPktCapturingState(true, &gpuPackets_, kernArgMgr,
&capturedKernelName_);
// 提交command以捕获AQL包 (不实际提交到设备)
command->submit(*(command->queue())->vdev());
command->release();
}
commands_.clear();
return status;
}
// 入队节点命令
virtual void EnqueueCommands(hip::Stream* stream) {
if (!isEnabled_) {
// 节点被禁用，变为空节点
amd::Command::EventWaitList waitList;
if (!commands_.empty()) {
waitList = commands_[0]->eventWaitList();
}
amd::Command* command = new amd::Marker(*stream, !kMarkerDisableFlush, waitList);
command->enqueue();
command->release();
return;
}
// 入队所有命令
for (auto& command : commands_) {
command->enqueue();
command->release();
}
}
protected:
hipGraphNodeType type_;
std::vector<Node> dependencies_; // 依赖节点
std::vector<Node> edges_; // 子节点
std::vector<uint8_t*> gpuPackets_; // 捕获的AQL包
std::vector<amd::Command*> commands_;
bool isEnabled_; // 节点是否启用
};
6.3 Graph捕获流程
文件: clr/hipamd/src/hip_graph.cpp:211-237
hipError_t capturehipLaunchKernel(
hipStream_t& stream,
const void*& hostFunction,
dim3& gridDim,
dim3& blockDim,
void**& args,
size_t& sharedMemBytes) {
hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
// 构建kernel节点参数
hipKernelNodeParams nodeParams;
nodeParams.func = const_cast<void*>(hostFunction);
nodeParams.blockDim = blockDim;
nodeParams.extra = nullptr;
nodeParams.gridDim = gridDim;
nodeParams.kernelParams = args;
nodeParams.sharedMemBytes = sharedMemBytes;
hip::GraphNode* pGraphNode;
// 添加kernel节点到捕获的graph
hipError_t status = ihipGraphAddKernelNode(
&pGraphNode, // 输出节点
s->GetCaptureGraph(), // 目标graph
s->GetLastCapturedNodes().data(), // 依赖节点
s->GetLastCapturedNodes().size(), // 依赖数量
&nodeParams, // 节点参数
nullptr, // 额外参数
true, // 捕获模式
0, // 标志
s->DeviceId() // 设备ID
);
// 更新最后捕获的节点
s->SetLastCapturedNode(pGraphNode);
return status;
}
6.4 AQL包捕获
文件: clr/rocclr/device/rocm/rocvirtual.cpp
Graph捕获模式下，dispatchAqlPacket 不会真正提交到GPU，而是将构建好的AQL包复制到Graph节点的缓冲区中。
bool VirtualGPU::dispatchAqlPacket(hsa_kernel_dispatch_packet_t* packet,
uint16_t header,
uint16_t rest,
bool blocking,
bool capturing,
const uint8_t* aqlPacket) {
if (capturing == true) {
// Graph捕获模式：将AQL包复制到Graph节点的缓冲区
packet->header = header;
packet->setup = rest;
std::memcpy(const_cast<uint8_t*>(aqlPacket), packet,
sizeof(hsa_kernel_dispatch_packet_t));
return true;
} else {
// 正常执行模式
dispatchBlockingWait();
return dispatchGenericAqlPacket(packet, header, rest, blocking);
}
}
捕获流程说明:
- capturing=true: Graph正在捕获阶段，AQL包被序列化到内存缓冲区
- capturing=false: 正常kernel执行，AQL包被提交到GPU队列
- aqlPacket参数: 指向Graph节点内部的
gpuPackets_缓冲区，用于存储捕获的AQL包
6.5 Graph捕获与生成流程
Graph捕获阶段
Graph Capture Phase (Stream Capture Mode)
═══════════════════════════════════════════
1. hipStreamBeginCapture(stream, mode)
└─▶ Stream进入CAPTURE模式
└─▶ 创建新的hipGraph对象
└─▶ 初始化捕获状态
2. 执行CUDA/HIP API (在捕获stream上)
├─▶ hipLaunchKernel()
│ └─▶ capturehipLaunchKernel()
│ └─▶ 创建GraphKernelNode
│ ├─▶ 记录kernel参数
│ ├─▶ 添加到Graph的节点列表
│ └─▶ 设置依赖关系
│
├─▶ hipMemcpyAsync()
│ └─▶ 创建GraphMemcpyNode
│
├─▶ hipMemsetAsync()
│ └─▶ 创建GraphMemsetNode
│
└─▶ ... (其他操作)
3. hipStreamEndCapture(stream, &graph)
└─▶ Stream退出CAPTURE模式
└─▶ 返回构建好的hipGraph对象
└─▶ Graph包含完整的节点拓扑结构
Graph实例化阶段
Graph Instantiate Phase (AQL Packet Generation)
════════════════════════════════════════════════
hipGraphInstantiate(graph, &graphExec)
│
├─▶ 1. 拓扑排序 (Topological Sort)
│ └─▶ 使用DFS/BFS算法排序节点
│ └─▶ 确保依赖节点先执行
│
├─▶ 2. 为每个节点生成AQL包
│ └─▶ GraphNode::CaptureAndFormPacket()
│ │
│ ├─▶ 创建临时Command
│ ├─▶ 设置捕获模式: setPktCapturingState(true, &gpuPackets_)
│ ├─▶ 提交Command到VirtualGPU (但不真正执行)
│ │ └─▶ VirtualGPU::submitKernelInternal()
│ │ └─▶ 构建hsa_kernel_dispatch_packet_t
│ │ ├─▶ kernel_object (kernel代码句柄)
│ │ ├─▶ kernarg_address (参数地址)
│ │ ├─▶ grid_size_* (global work size)
│ │ └─▶ workgroup_size_* (local work size)
│ │
│ └─▶ dispatchAqlPacket(packet, header, rest, blocking,
│ /*capturing=*/true, aqlPacket)
│ └─▶ 将AQL包复制到gpuPackets_缓冲区
│
└─▶ 3. 创建hipGraphExec对象
└─▶ 包含所有预生成的AQL包
└─▶ 包含节点依赖关系
└─▶ 可用于多次重放
Graph执行阶段
Graph Launch Phase (AQL Packet Replay)
═══════════════════════════════════════
hipGraphLaunch(graphExec, stream)
│
├─▶ 1. 遍历拓扑排序的节点
│ └─▶ for node in sorted_nodes:
│ └─▶ node->EnqueueCommands(stream)
│
├─▶ 2. 提交预生成的AQL包
│ └─▶ 从gpuPackets_获取AQL包
│ └─▶ 复制到GPU队列内存
│ └─▶ dispatchGenericAqlPacket()
│ ├─▶ 写入队列缓冲区
│ └─▶ 触发Doorbell
│
└─▶ 3. GPU执行
└─▶ CP读取AQL包
└─▶ 按依赖顺序执行kernel
└─▶ 更新completion_signal
Graph优势对比
传统Kernel Launch vs Graph Launch
═══════════════════════════════════
传统方式 (每次执行):
┌─────────────────────────────────────────────────────────┐
│ 1. 构建kernel参数 │
│ 2. 创建NDRangeKernelCommand │
│ 3. 构建AQL packet (64字节) │
│ 4. 获取队列写索引 (原子操作) │
│ 5. 写入AQL packet到队列内存 │
│ 6. 触发Doorbell (MMIO写操作) │
└─────────────────────────────────────────────────────────┘
重复以上步骤N次
Graph方式 (首次实例化后):
┌─────────────────────────────────────────────────────────┐
│ 实例化阶段 (一次性): │
│ - 拓扑排序 │
│ - 预生成所有AQL包 │
│ - 存储在GraphExec中 │
└─────────────────────────────────────────────────────────┘
每次执行只需:
┌─────────────────────────────────────────────────────────┐
│ 1. 复制预生成的AQL包到队列 (memcpy) │
│ 2. 触发Doorbell │
└─────────────────────────────────────────────────────────┘
性能提升:
- 减少CPU开销: 避免重复的AQL包构建
- 减少系统调用: 预计算所有参数
- 更好的调度: 静态依赖分析优化
- 适合重复执行: 推理场景、循环计算
完整调用流程图
7.1 从MIOpen到GPU的完整流程
┌──────────────────────────────────────────────────────────────────────────────┐
│ MIOpen层 (用户API) │
├──────────────────────────────────────────────────────────────────────────────┤
│ │
│ miopenPoolingForward() │
│ └─▶ pooling_api.cpp:312 │
│ └─▶ PoolingDescriptor::Forward() │
│ └─▶ pooling_ocl.cpp:57 │
│ └─▶ PoolingForwardSolvers().ExecutePrimitive() │
│ └─▶ Solver执行kernel │
│ │
├──────────────────────────────────────────────────────────────────────────────┤
│ MIOpen Kernel层 │
├──────────────────────────────────────────────────────────────────────────────┤
│ │
│ Handle::Run() │
│ └─▶ handlehip.cpp:527 │
│ └─▶ HIPOCKernelInvoke::run() │
│ └─▶ hipoc_kernel.cpp:86 │
│ └─▶ hipExtModuleLaunchKernel() │
│ │
├──────────────────────────────────────────────────────────────────────────────┤
│ HIP Runtime层 │
├──────────────────────────────────────────────────────────────────────────────┤
│ │
│ hipExtModuleLaunchKernel() │
│ └─▶ hip_module.cpp:468 │
│ └─▶ ihipModuleLaunchKernel() │
│ └─▶ hip_module.cpp:357 │
│ └─▶ 创建 amd::NDRangeKernelCommand │
│ └─▶ kernelCommand->enqueue() │
│ │
├──────────────────────────────────────────────────────────────────────────────┤
│ ROCclr层 (命令队列) │
├──────────────────────────────────────────────────────────────────────────────┤
│ │
│ Command::enqueue() │
│ └─▶ command.cpp │
│ └─▶ queue_->FormSubmissionBatch(this) │
│ └─▶ VirtualGPU::submitKernel() │
│ └─▶ rocvirtual.cpp:3414 │
│ └─▶ submitKernelInternal() │
│ └─▶ rocvirtual.cpp:2993 │
│ ┌─────────────────────────────────────┐ │
│ │ 1. 创建 hsa_kernel_dispatch_packet_t │ │
│ │ 2. 设置 kernel_object │ │
│ │ 3. 设置 grid/workgroup 大小 │ │
│ │ 4. 设置 kernarg_address │ │
│ │ 5. 调用 dispatchAqlPacket() │ │
│ └─────────────────────────────────────┘ │
│ │
├──────────────────────────────────────────────────────────────────────────────┤
│ AQL包分发层 │
├──────────────────────────────────────────────────────────────────────────────┤
│ │
│ VirtualGPU::dispatchGenericAqlPacket() │
│ └─▶ rocvirtual.cpp:840 │
│ ├─▶ hsa_queue_add_write_index_screlease() // 获取写索引 │
│ ├─▶ Barriers().ActiveSignal() // 获取完成信号 │
│ ├─▶ 写入AQL包到 gpu_queue_->base_address // 队列内存 │
│ ├─▶ packet_store_release() // 写入header │
│ └─▶ hsa_signal_store_screlease() // 触发Doorbell │
│ │
├──────────────────────────────────────────────────────────────────────────────┤
│ ROCR Runtime层 (HSA) │
├──────────────────────────────────────────────────────────────────────────────┤
│ │
│ hsa_signal_store_screlease() │
│ └─▶ hsa.cpp:1209 │
│ └─▶ AqlQueue::StoreRelaxed() │
│ └─▶ amd_aql_queue.cpp:461 │
│ └─▶ 写入 doorbell指针 (MMIO) │
│ │
├──────────────────────────────────────────────────────────────────────────────┤
│ ROCT-Thunk层 │
├──────────────────────────────────────────────────────────────────────────────┤
│ │
│ hsaKmtCreateQueue() / hsaKmtUpdateQueue() │
│ └─▶ libhsakmt/src/queues.c │
│ └─▶ hsakmt_ioctl() │
│ └─▶ ioctl() 系统调用 │
│ │
├──────────────────────────────────────────────────────────────────────────────┤
│ KFD Kernel Driver层 │
├──────────────────────────────────────────────────────────────────────────────┤
│ │
│ kfd_ioctl_create_queue() / kfd_ioctl_update_queue() │
│ └─▶ amdkfd/kfd_chardev.c │
│ └─▶ 配置GPU CP (Command Processor) │
│ └─▶ 设置队列基地址和Doorbell │
│ │
├──────────────────────────────────────────────────────────────────────────────┤
│ GPU Hardware层 │
├──────────────────────────────────────────────────────────────────────────────┤
│ │
│ Command Processor (CP) │
│ ├─▶ 检测Doorbell变化 │
│ ├─▶ 从队列读取AQL Packet │
│ ├─▶ 解析Packet Header │
│ ├─▶ 分配Shader Engine │
│ ├─▶ 加载Kernel代码 │
│ ├─▶ 设置寄存器 │
│ └─▶ 启动Wavefronts执行 │
│ │
│ Shader CUs (Compute Units) │
│ └─▶ 执行Kernel指令 │
│ └─▶ 完成时更新 completion_signal │
│ │
└──────────────────────────────────────────────────────────────────────────────┘
7.2 AQL Packet内存布局
AQL Kernel Dispatch Packet (64字节对齐)
Offset Field Size Description
─────────────────────────────────────────────────────────────
0x00 header 2 bytes Packet type, barrier, fence
0x02 setup 2 bytes Dimensions (1D/2D/3D)
0x04 workgroup_size_x 2 bytes Workgroup X dimension
0x06 workgroup_size_y 2 bytes Workgroup Y dimension
0x08 workgroup_size_z 2 bytes Workgroup Z dimension
0x0A reserved0 2 bytes Reserved
0x0C grid_size_x 4 bytes Grid X dimension
0x10 grid_size_y 4 bytes Grid Y dimension
0x14 grid_size_z 4 bytes Grid Z dimension
0x18 private_segment_size 4 bytes Private memory size
0x1C group_segment_size 4 bytes LDS (Local Data Share) size
0x20 kernel_object 8 bytes Kernel code object handle
0x28 kernarg_address 8 bytes Kernel arguments address
0x30 reserved2 8 bytes Reserved
0x38 completion_signal 8 bytes Completion signal handle
─────────────────────────────────────────────────────────────
Total: 64 bytes
Header Format (16 bits):
Bits 0-7: Type (1=Kernel, 2=Barrier-AND, 3=Barrier-OR)
Bit 8: Barrier (1=wait for completion)
Bits 9-10: Acquire fence scope
Bits 11-12: Release fence scope
Bits 13-15: Reserved
7.3 Graph执行流程图
Graph Capture Phase Graph Execution Phase
───────────────────── ─────────────────────
hipStreamBeginCapture() hipGraphInstantiate()
│ │
▼ ▼
┌─────────────┐ ┌─────────────────┐
│ Stream进入 │ │ 拓扑排序节点 │
│ CAPTURE模式 │ │ (DFS/BFS) │
└──────┬──────┘ └────────┬────────┘
│ │
▼ ▼
hipLaunchKernel() ┌─────────────┐
│ │ 为每个节点 │
▼ │ 调用Capture │
┌─────────────┐ │ AndFormPacket│
│ 创建GraphNode│ │ │
│ (kernel类型) │ └──────┬──────┘
└──────┬──────┘ │
│ ▼
▼ ┌─────────────┐
hipMemcpyAsync() │ 生成AQL包 │
│ │ 存储在节点 │
▼ │ gpuPackets_ │
┌─────────────┐ └──────┬──────┘
│ 创建GraphNode│ │
│ (memcpy类型) │ ▼
└──────┬──────┘ ┌─────────────┐
│ │ 创建 │
▼ │ GraphExec │
... (更多操作) │ 对象 │
└─────────────┘
│
▼
hipStreamEndCapture()
│
▼
┌─────────────┐
│ 得到hipGraph │
│ 对象 │
└─────────────┘
hipGraphLaunch(graphExec, stream)
│
▼
┌─────────────────────────────────────────┐
│ 遍历拓扑排序的节点 │
│ for each node in sorted_nodes: │
│ node->EnqueueCommands(stream) │
└─────────────────────────────────────────┘
│
▼
┌─────────────────────────────────────────┐
│ 提交预生成的AQL包 │
│ - 复制AQL包到队列内存 │
│ - 触发Doorbell │
└─────────────────────────────────────────┘
│
▼
GPU执行 (同普通kernel执行)
关键文件路径汇总
| 组件 | 文件路径 | 功能描述 |
|---|---|---|
| MIOpen | ||
| Kernel Cache | MIOpen/src/kernel_cache.cpp |
Kernel编译缓存管理 |
| HIP Program | MIOpen/src/hipoc/hipoc_program.cpp |
HIP程序编译和加载 |
| HIP Kernel | MIOpen/src/hipoc/hipoc_kernel.cpp |
Kernel执行封装 |
| Handle HIP | MIOpen/src/hip/handlehip.cpp |
HIP设备管理 |
| Pooling API | MIOpen/src/pooling_api.cpp |
Pooling操作API |
| HIP Runtime | ||
| Module | clr/hipamd/src/hip_module.cpp |
Module加载和kernel启动 |
| Graph | clr/hipamd/src/hip_graph.cpp |
Graph API实现 |
| Graph Internal | clr/hipamd/src/hip_graph_internal.hpp |
Graph节点定义 |
| ROCclr | ||
| Command Queue | clr/rocclr/platform/commandqueue.cpp |
命令队列管理 |
| Command | clr/rocclr/platform/command.cpp |
命令基类实现 |
| VirtualGPU | clr/rocclr/device/rocm/rocvirtual.cpp |
AQL包生成和提交 |
| ROCR Runtime | ||
| HSA Header | ROCR-Runtime/runtime/hsa-runtime/inc/hsa.h |
AQL结构定义 |
| AQL Queue | ROCR-Runtime/runtime/hsa-runtime/core/runtime/amd_aql_queue.cpp |
Queue实现 |
| ROCT-Thunk | ||
| Queues | ROCR-Runtime/libhsakmt/src/queues.c |
Queue创建和管理 |
| KFD Driver | ||
| Chardev | ROCK-Kernel-Driver/drivers/gpu/drm/amd/amdkfd/kfd_chardev.c |
KFD字符设备 |
总结
本文档详细分析了ROCm软件栈中kernel执行的完整流程：
-
Kernel生成: MIOpen通过Kernel Cache管理编译后的kernel，使用COMgr编译HIP代码，加载为HIP module。
-
Kernel入队: HIP Runtime创建NDRangeKernelCommand，通过ROCclr的命令队列管理系统入队。
-
AQL包: ROCclr生成标准的HSA AQL packet，包含kernel执行所需的全部信息（代码地址、参数、维度等）。
-
队列管理: AQL包被写入GPU队列的环形缓冲区，通过Doorbell机制通知GPU。
-
Graph API: 预先生成AQL包并存储在Graph节点中，执行时直接提交，减少CPU开销。
-
硬件执行: GPU的Command Processor检测Doorbell，读取AQL包，调度Shader CUs执行kernel。
理解这个流程对于优化ROCm应用程序性能、调试问题和开发新功能都非常重要。
附录: 完整Kernel编译流程图
┌─────────────────────────────────────────────────────────────────────────────┐
│ MIOpen Kernel Compilation Flow │
└─────────────────────────────────────────────────────────────────────────────┘
1. API CALL
miopenPoolingForward()
│
v
2. PROBLEM DESCRIPTION
PoolingDescriptor::Forward()
└── Creates: ProblemDescription, FwdInvokeParams
│
v
3. SOLVER SELECTION
SolverContainer::ExecutePrimitive()
└── SearchForSolutions() → Finds first applicable solver
│
v
4. SOLUTION GENERATION
Solver::GetSolution()
└── Returns: ConvSolution with:
├── KernelInfo (kernel_file, kernel_name, comp_options, g_wk, l_wk)
└── invoker_factory (lambda to create kernel invoker)
│
v
5. INVOKER PREPARATION
Handle::PrepareInvoker()
└── KernelCache::AddKernel()
│
v
6. PROGRAM LOADING/COMPILATION
Handle::LoadProgram()
├── 1. Check Binary Cache (LoadBinary)
│ └── Found? Return cached binary
│
└── 2. Compile from Source (Cache Miss)
│
├── HIPOCProgram Constructor
│ └── BuildCodeObject()
│ ├── GetKernelSrc() - Get embedded source
│ │
│ └── #if MIOPEN_USE_COMGR
│ ├── .cpp → hiprtc::BuildHip()
│ ├── .s → comgr::BuildAsm()
│ └── .cl → comgr::BuildOcl()
│ │
│ └── COMgr Actions:
│ ├── AMD_COMGR_ACTION_ADD_PRECOMPILED_HEADERS
│ ├── AMD_COMGR_ACTION_COMPILE_SOURCE_WITH_DEVICE_LIBS_TO_BC
│ ├── AMD_COMGR_ACTION_CODEGEN_BC_TO_RELOCATABLE
│ └── AMD_COMGR_ACTION_LINK_RELOCATABLE_TO_EXECUTABLE
│
└── 3. Save to Binary Cache (SaveBinary)
│
v
7. KERNEL CREATION
HIPOCKernel Constructor
└── hipModuleGetFunction() - Extract kernel from module
│
v
8. KERNEL EXECUTION
HIPOCKernelInvoke::operator()
└── hipModuleLaunchKernel() - Execute on GPU
关键文件路径汇总
| 组件 | 文件路径 | 功能描述 |
|---|---|---|
| MIOpen API | ||
| Pooling API | MIOpen/src/pooling_api.cpp |
API入口 |
| Pooling Implementation | MIOpen/src/ocl/pooling_ocl.cpp |
Pooling实现 |
| Solver系统 | ||
| Solver Base | MIOpen/src/include/miopen/solver.hpp |
Solver基类 |
| Solver Selection | MIOpen/src/include/miopen/find_solution.hpp |
Solution查找 |
| Pooling Solvers | MIOpen/src/include/miopen/pooling/solvers.hpp |
Pooling solver定义 |
| PoolingForward2d | MIOpen/src/solver/pooling/forward2d.cpp |
2D pooling solver |
| 编译系统 | ||
| HIP Handle | MIOpen/src/hip/handlehip.cpp |
Handle实现 |
| Kernel Cache | MIOpen/src/kernel_cache.cpp |
Kernel缓存管理 |
| HIPOC Program | MIOpen/src/hipoc/hipoc_program.cpp |
Program编译 |
| COMgr Integration | MIOpen/src/comgr.cpp |
COMgr编译器接口 |
| Binary Cache | MIOpen/src/binary_cache.cpp |
Binary缓存系统 |
| Kernel源文件 | ||
| Kernel Sources | MIOpen/src/kernels/MIOpenPooling.cl |
Pooling kernel源码 |
| Kernel Embed | MIOpen/src/kernels/kernel.cpp.in |
源码嵌入模板 |
| 执行系统 | ||
| HIPOC Kernel | MIOpen/src/include/miopen/hipoc_kernel.hpp |
Kernel封装 |
免费领 150 小时云算力，进群参与显卡、AI PC 幸运抽奖
更多推荐
11
0- 0
已为社区贡献2条内容
所有评论(0)