Skip to content
Merged
Show file tree
Hide file tree
Changes from 33 commits
Commits
Show all changes
44 commits
Select commit Hold shift + click to select a range
3fd78d1
(CMCL) hip-kernel-provider: add the saved ingestor plan payload format
CMiservaAMD Oct 2, 2026
a7dc5ca
(CMCL) hip-kernel-provider: add HIPKERNELPROVIDER_GENERATE_SERIALIZE_…
CMiservaAMD Oct 4, 2026
a7d66d1
(CMCL) hip-kernel-provider: refine the saved ingestor plan payload sc…
CMiservaAMD Oct 5, 2026
0cd9953
(CMCL) hip-kernel-provider: capture an ingestor plan for saving
CMiservaAMD Oct 5, 2026
a2c2354
(CMCL) hip-kernel-provider: restore a saved ingestor plan
CMiservaAMD Oct 7, 2026
732688a
(CMCL) hip-kernel-provider: export the plan serialization hooks
CMiservaAMD Oct 8, 2026
d70ddca
(CMCL) hip-kernel-provider: drop engine identity from saved plans
CMiservaAMD Oct 8, 2026
cfb0b3e
(CMCL) hip-kernel-provider: test and document saved ingestor plans
CMiservaAMD Oct 8, 2026
8807207
(CMCL) hip-kernel-provider: consolidate the saved-plan tests
CMiservaAMD Oct 8, 2026
2c325f2
(CMCL) hip-kernel-provider: restore scale check and fast-path note test
CMiservaAMD Oct 9, 2026
6a7e858
(CMCL) hip-kernel-provider: read kpack code bytes through one reader
CMiservaAMD Oct 9, 2026
838e6d6
(CMCL) hip-kernel-provider: hand the saved buffer over last
CMiservaAMD Oct 13, 2026
f3b865a
(CMCL) hip-kernel-provider: narrow the dispatch capability probe
CMiservaAMD Oct 13, 2026
fc5f786
(CMCL) hip-kernel-provider: give the plan schema explicit field ids
CMiservaAMD Oct 13, 2026
2be6797
(CMCL) hip-kernel-provider: fail loudly when the plan header is not r…
CMiservaAMD Oct 13, 2026
257e41a
(CMCL) hipdnn: resolve engine names through the non-throwing lookup
CMiservaAMD Oct 13, 2026
49d267a
(CMCL) hip-kernel-provider: sweep every resealed plan body byte
CMiservaAMD Oct 13, 2026
1730568
(CMCL) hip-kernel-provider: pin the launch value contracts with saved…
CMiservaAMD Oct 13, 2026
e0a3184
(CMCL) hip-kernel-provider: correct the saved-sample and regeneration…
CMiservaAMD Oct 13, 2026
d160a44
[hipDNN] Document the memory contract of the plugin buffer hooks (ALM…
CMiservaAMD Oct 14, 2026
dfed5e4
(CMCL) hip-kernel-provider: own the ingestor plan flatc flags and def…
CMiservaAMD Oct 16, 2026
b19afb3
(CMCL) hip-kernel-provider: pin the ingestor plan schema hash beside …
CMiservaAMD Oct 16, 2026
54d4996
(CMCL) hip-kernel-provider: refuse a saved plan whose kernel digest d…
CMiservaAMD Oct 16, 2026
041ceca
(CMCL) hip-kernel-provider: clang-format the kernel digest refusal test
CMiservaAMD Oct 16, 2026
2c5b083
(CMCL) hip-kernel-provider: encode a saved plan straight into the buf…
CMiservaAMD Oct 16, 2026
6d7eeb8
(CMCL) hip-kernel-provider: compare the whole saved payload in the la…
CMiservaAMD Oct 16, 2026
a2d767c
(CMCL) hip-kernel-provider: pin the conv and pointwise launch geometr…
CMiservaAMD Oct 16, 2026
e17d7a0
(CMCL) hip-kernel-provider: say a save may fail when the kpack archiv…
CMiservaAMD Oct 16, 2026
b194537
(CMCL) hip-kernel-provider: accept a null buffer in the serialized-pl…
CMiservaAMD Oct 16, 2026
0ea9a72
(CMCL) hip-kernel-provider: test that a restored kernel loads a new d…
CMiservaAMD Oct 16, 2026
9dec3ca
(CMCL) hipdnn: lock the serialized execution plan cache (ALMIOPEN-2941)
CMiservaAMD Oct 16, 2026
0eeefd0
(CMCL) hip-kernel-provider: declare the null-bytes test structure const
CMiservaAMD Oct 16, 2026
93831e5
(CMCL) hip-kernel-provider: scope the schema change bans to a minor v…
CMiservaAMD Oct 16, 2026
bba2d23
(CMCL) hipdnn: fail cleanly when the retried plan save returns the wr…
CMiservaAMD Oct 19, 2026
c33e3c9
(CMCL) hip-kernel-provider: print the signature and dispatch symbol w…
CMiservaAMD Oct 19, 2026
b5c955b
Merge dev/develop (299f298eabe) into almiopen-2642/save-execution-plans
CMiservaAMD Oct 19, 2026
91e9fd9
(CMCL) hip-kernel-provider: record the GPU a plan was saved on, also …
CMiservaAMD Oct 19, 2026
b33acfa
(CMCL) hipdnn plugin SDK, hip-kernel-provider: report plan serializat…
CMiservaAMD Oct 19, 2026
3d85e09
Merge develop (061b772a872) into almiopen-2642/save-execution-plans
CMiservaAMD Oct 20, 2026
95e79cb
Merge branch 'develop' into users/cmiserva/almiopen-2642/save-executi…
CMiservaAMD Oct 21, 2026
4d2fb10
Merge branch 'develop' into users/cmiserva/almiopen-2642/save-executi…
CMiservaAMD Oct 21, 2026
573f647
Merge develop (0bb1c36e8f7) into almiopen-2642/save-execution-plans
CMiservaAMD Oct 22, 2026
6e80ed1
Merge develop (9288d8b78fd) into almiopen-2642/save-execution-plans
CMiservaAMD Oct 22, 2026
3b93962
(CMCL) hipdnn docs: re-verify RFC 0009 against the plan serialization…
CMiservaAMD Oct 22, 2026
File filter

Filter by extension

Filter by extension


Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
11 changes: 11 additions & 0 deletions .github/workflows/hipdnn-superbuild-ci.yml
Original file line number Diff line number Diff line change
Expand Up @@ -220,6 +220,7 @@ jobs:
-DHIPKERNELPROVIDER_ENABLE_ROCKE=ON \
-DHIPDNN_ENABLE_CUDNN_COMPATIBILITY=ON \
-DHIPDNN_ENABLE_KERNEL_INGESTOR=ON \
-DHIPKERNELPROVIDER_GENERATE_SERIALIZE_HEADERS=ON \
-DHIPDNN_ENABLE_SDPA=ON \
-DHIPKERNELPROVIDER_KPACK_PYTHON_DIR="$KPACK_PYTHON_DIR" \
-DENABLE_CLANG_TIDY=OFF \
Expand Down Expand Up @@ -462,6 +463,7 @@ jobs:
-DHIPKERNELPROVIDER_ENABLE_ROCKE=ON \
-DHIPDNN_ENABLE_CUDNN_COMPATIBILITY=ON \
-DHIPDNN_ENABLE_KERNEL_INGESTOR=ON \
-DHIPKERNELPROVIDER_GENERATE_SERIALIZE_HEADERS=ON \
-DHIPDNN_ENABLE_SDPA=ON \
-DHIPKERNELPROVIDER_PRODUCTION_SOURCE_ROOT="$GITHUB_WORKSPACE/dnn-providers/hip-kernel-provider/descriptor-packaging/examples/descriptors" \
-DHIPKERNELPROVIDER_ROCKE_COMGR_LIB="$HKP_ROCKE_COMGR_LIB" \
Expand All @@ -475,6 +477,15 @@ jobs:
- name: Build
run: cmake --build build

# The touch makes the build run flatc again, also when the checkout gives the committed
# header a newer timestamp than the schema.
- name: Check the committed ingestor plan header matches its schema
run: |
set -euo pipefail
touch dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/serialization/ingestor_plan.fbs
cmake --build build --target hkp_generate_ingestor_plan_header
git diff --exit-code -- dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/serialization/ingestor_plan_generated.h

# Packing is skipped with a STATUS message, so a resolution regression
# would otherwise pass CI while producing no archives.
- name: Verify kpack archives were produced
Expand Down
2 changes: 2 additions & 0 deletions .pre-commit-config.yaml
Original file line number Diff line number Diff line change
Expand Up @@ -49,6 +49,8 @@ exclude: |
projects/hipdnn/.*\.bin$|
# Generated hipDNN FlatBuffers headers
projects/hipdnn/flatbuffers_sdk/include/hipdnn_flatbuffers_sdk/data_objects/.*_generated\.h|
# Generated hip-kernel-provider FlatBuffers header
dnn-providers/hip-kernel-provider/src/engines/kernel_ingestor_engine/serialization/ingestor_plan_generated\.h|
# Reference data (auto-generated)
dnn-providers/integration-tests/integration-test-bundles/.*|
# Generated files (see rocrand/tools/ _generator.cpp files)
Expand Down
3 changes: 3 additions & 0 deletions dnn-providers/hip-kernel-provider/.clang-format-ignore
Original file line number Diff line number Diff line change
Expand Up @@ -3,3 +3,6 @@
# templates for the kernel inlining -- clang-format breaks the
# syntax of varialbes by adding a space after the second @ character
src/cmake/templates/*

# FlatBuffers output, generated from ingestor_plan.fbs by flatc
src/engines/kernel_ingestor_engine/serialization/ingestor_plan_generated.h
28 changes: 23 additions & 5 deletions dnn-providers/hip-kernel-provider/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -47,6 +47,16 @@ list(APPEND CMAKE_MODULE_PATH ${CMAKE_CURRENT_SOURCE_DIR}/cmake

option(HIPKERNELPROVIDER_ENABLE_TESTS "Build tests" ON)
option(HIPKERNELPROVIDER_ENABLE_ROCKE "Build rocKE engine and smoke tests" OFF)
# A TheRock build compiles the committed serialization headers by default.
if(DEFINED THEROCK_SUBPROJECT_TARGET)
set(_hkp_generate_serialize_default OFF)
else()
set(_hkp_generate_serialize_default ON)
endif()
option(HIPKERNELPROVIDER_GENERATE_SERIALIZE_HEADERS
"Regenerate FlatBuffers serialization headers from schema during build"
${_hkp_generate_serialize_default}
)

# Gates the generic kernel ingestor. The ingestor engine compiles its kernel
# through the HIP kernel engine's compiler, so it only builds alongside that engine.
Expand Down Expand Up @@ -311,7 +321,9 @@ if(HIPDNN_ENABLE_KERNEL_INGESTOR)
)

# Two test descriptor roots, one per test binary. The unit root holds two packed
# subtrees and the integration root holds three, all five wired in
# subtrees and the integration root holds three. A sixth packed subtree, the unit
# binary's copy of archive_fixture, stages beside the two roots and not inside the unit
# root, so the unit engine does not discover it. All six are wired in
# descriptor-packaging/cmake/HkpPackaging.cmake. The packer stages each authored
# subtree per architecture: it rewrites an authored hip descriptor to kpack and writes
# the archive beside it. An embedded_source descriptor names a kernel compiled into a
Expand All @@ -329,8 +341,8 @@ if(HIPDNN_ENABLE_KERNEL_INGESTOR)

# Not a test root: the production tree's provenance sidecars install here. It shares
# the folder with the test sets, so it must not reuse unit, integration,
# test_kpack_binary (src/tests/CMakeLists.txt) or the census folder HkpPackaging.cmake
# installs.
# unit_archive_fixture, test_kpack_binary (src/tests/CMakeLists.txt) or the census
# folder HkpPackaging.cmake installs.
set(HIPKERNELPROVIDER_TEST_SET_PROVENANCE "provenance")

set(HIPKERNELPROVIDER_TEST_DESCRIPTOR_SOURCE_ROOT
Expand All @@ -343,8 +355,8 @@ if(HIPDNN_ENABLE_KERNEL_INGESTOR)

# The name of each authored set, which is also the name it stages under. Named here
# only because the test subdirectories compile these same segments into their binaries
# as a relative offset; the packing rules in HkpPackaging.cmake join them to the two
# roots above directly.
# as a relative offset; the packing rules in HkpPackaging.cmake join them to the roots
# above.
#
# Every set is a top-level folder under test_descriptors/, and every one is packed
# WHOLE rather than by the set inside it. Each descriptor therefore lands in a CHILD of
Expand All @@ -361,6 +373,12 @@ if(HIPDNN_ENABLE_KERNEL_INGESTOR)
set(HIPKERNELPROVIDER_TEST_SET_SHARED "shared")
set(HIPKERNELPROVIDER_TEST_SET_ARCHIVE_FIXTURE "archive_fixture")

# The unit binary's own packed copy of the archive_fixture set. The integration suite
# corrupts the integration copy in place, so the two binaries do not share an archive.
set(HIPKERNELPROVIDER_UNIT_ARCHIVE_FIXTURE_SUBDIR "unit_archive_fixture")
set(HIPKERNELPROVIDER_UNIT_ARCHIVE_FIXTURE_BUILD_DIR
"${HIPKERNELPROVIDER_TEST_DESCRIPTOR_BUILD_DIR}/${HIPKERNELPROVIDER_UNIT_ARCHIVE_FIXTURE_SUBDIR}")

# The same binary-to-plugin-directory offset, extended to the test descriptor trees.
file(RELATIVE_PATH HIPKERNELPROVIDER_TEST_DESCRIPTOR_RELDIR
"${_synthetic_root}/${CMAKE_INSTALL_BINDIR}"
Expand Down
85 changes: 85 additions & 0 deletions dnn-providers/hip-kernel-provider/README.md
Original file line number Diff line number Diff line change
Expand Up @@ -96,6 +96,91 @@ Kernel source files (`.cpp`, `.hpp`, `.h`) under `engines/` are embedded as C++

The embedding is handled by the `embed_kernel_sources()` CMake function in `src/cmake/KernelEmbedding.cmake`.

## Saving and Loading Execution Plans

You can save a built execution plan to bytes and load it later, in the same process or in another one. A loaded plan runs without the plan build: the provider does not search descriptors, read a kpack archive, or compile a kernel.

### Supported Engines

Only kernel ingestor engines support saving and loading plans, and only when every kernel of the engine is a kpack code object. The kernel ingestor needs a build with `HIPDNN_ENABLE_KERNEL_INGESTOR=ON` (off by default). A supported engine reports the behavior note `SUPPORTS_EXECUTION_PLAN_SERIALIZATION`. To check the engine of a built plan:

```cpp
#include <algorithm>
#include <hipdnn_frontend.hpp>

int64_t engineId = 0;
std::vector<hipdnn_frontend::BehaviorNote> notes;
bool canSave = false;
if(graph.get_execution_plan_engine_id(engineId).is_good()
&& graph.get_behavior_notes_for_engine(engineId, notes).is_good())
{
canSave = std::find(notes.begin(),
notes.end(),
hipdnn_frontend::BehaviorNote::SUPPORTS_EXECUTION_PLAN_SERIALIZATION)
!= notes.end();
}
if(canSave)
{
// Save the plan with to_binary() or to_compiled_plan_binary().
}
```

The other engines of this provider do not report the note.

### Saving the Graph and the Plan Together

`to_binary()` writes the graph. When the engine of the built plan reports the note, it also writes the plan. `from_binary(handle, bytes)` loads both, and the loaded graph executes without a new build.

```cpp
// graph is built (for example with graph.build(handle)) and has executed once.
auto [bytes, error] = graph.to_binary();
if(error.is_bad())
{
std::cerr << error.get_message() << '\n';
}

hipdnn_frontend::graph::Graph loaded;
auto loadError = loaded.from_binary(handle, bytes);
// loaded.execute(handle, variantPack, workspace);
```

When the engine does not report the note, `to_binary()` writes the graph only and logs a warning. A graph loaded from those bytes has no plan, so build its plans again before you execute it.

### Saving the Plan Only

`to_compiled_plan_binary()` writes the plan only. `from_compiled_plan_binary(handle, bytes)` loads it into a new graph. The loaded graph has no operation graph; execute it with a variant pack keyed by tensor UID.

```cpp
auto [planBytes, error] = graph.to_compiled_plan_binary();
if(error.is_bad())
{
std::cerr << error.get_message() << '\n';
}

hipdnn_frontend::graph::Graph restored;
auto loadError = restored.from_compiled_plan_binary(handle, planBytes);

std::unordered_map<int64_t, void*> variantPack = {{xUid, xDevicePtr}, {yUid, yDevicePtr}};
auto executeError = restored.execute(handle, variantPack, workspace);
```

A plan-only save of a plan from any other engine of this provider fails. With `HIPDNN_ENABLE_KERNEL_INGESTOR=ON` the backend status is `HIPDNN_STATUS_PLUGIN_ERROR`; without the ingestor the provider has no save support and the status is `HIPDNN_STATUS_NOT_SUPPORTED`. In both cases the frontend returns `ErrorCode::HIPDNN_BACKEND_ERROR`, with the cause in the message.

### Rules

- **Benchmarking plans save only after they have chosen a kernel.** A plan built with `global.benchmarking=1` or with `HIPDNN_FORCE_BENCHMARKING=1` usually chooses its kernel at its first `execute()`. (When a stored benchmark result already covers every candidate kernel, the build makes a plain plan that saves at once.) Saving a plan that has not chosen a kernel fails with "the plan has not chosen a kernel yet; execute it once, or build it with benchmarking off, then save". Execute the plan once and then save it, or build it with benchmarking off. For an engine that reports the note, this also makes `to_binary()` fail; it does not fall back to the graph only.
- **A loaded plan cannot be saved again.** Both `to_compiled_plan_binary()` of a loaded plan and `to_binary()` of a graph loaded with `from_binary(handle, bytes)` fail with "this plan was loaded from a saved payload; re-saving is not supported; keep the original bytes". Keep the bytes you loaded, and reuse them.
- **A loaded plan runs the code it was saved with.** Upgrading the provider does not change the kernel inside a saved plan, even when the upgrade fixes that kernel. To use the new kernel, build the plan again from its graph and save it.
- **The kpack archive must not change between the plan build and the save.** A save gets the kernel's code object from the kpack archive that the plan build used, and checks it against the kernel's SHA-256. The save can read the archive through the handle that the plan build opened. If the archive changes after the plan build, the save may fail. If the save cannot read the archive, or the bytes do not match, the save fails.
- **The loading machine must have the same engine.** hipDNN loads a plan only through a plugin that serves the saved engine ID, so the descriptor set that defines the engine must be installed and loaded. The GPU must be able to run the stored code object, and this provider must still recognize the stored dispatch name, kernel arguments and launch values. The format version must be one this provider reads. The provider version does not need to match.
- **Failures name their cause.** When the provider refuses a save or a load, the backend status is `HIPDNN_STATUS_PLUGIN_ERROR`. When no loaded plugin serves the saved engine ID at all, hipDNN refuses the load before the provider runs, with `HIPDNN_STATUS_INTERNAL_ERROR` and "Invalid engine ID". In every case the frontend returns `ErrorCode::HIPDNN_BACKEND_ERROR`. The error text holds hipDNN's own text first and then the provider's message. The provider's message starts with "ingestor plan data is damaged: " or "cannot save this plan: its kernel source does not match its descriptor: " when the data is damaged or the kernel no longer matches its descriptor, and with "ingestor plan is valid but cannot be saved or loaded here: " or "cannot save this plan here: " when the plan is valid but cannot be used here. Search the error text for these prefixes; do not expect it to start with them. A load checks everything, including that the GPU driver accepts the code object, so a bad plan fails at load and not at its first execute.

### Security

A saved plan contains executable GPU code, and executing a loaded plan runs that code. Load only plans from a source you trust. The SHA-256 checks detect accidental damage, not a crafted file. A load does not repeat the path checks the provider applies to its descriptor tree, and the code object goes to the GPU driver as it is.

The full format, the load checks and their order are in [RFC 0009](../../projects/hipdnn/docs/rfcs/0009_CompiledPlanSerialization.md#hip-kernel-provider-ingestor-implementation).

## Contributing

When adding new operations:
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -34,8 +34,9 @@ A root, or an arch, left with nothing to pack is skipped, not failed. Only produ
gated by a build option (rocKE today) are disabled by kind; a kind outside that set is
not pruned.

The provider wires six roots: production, plus five over the four authored test sets
(`shared` packs twice, once into each test binary's discovery root).
The provider wires seven roots: production, plus six over the four authored test sets.
`shared` packs twice, once into each test binary's discovery root. `archive_fixture` packs
twice, once for each test binary, so each binary reads its own archive.

| Root | Source | Ships |
|---|---|---|
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -1653,7 +1653,7 @@ function(hkp_add_packaging)

# Test descriptors, one pack per authored set. The shared root is packed into both
# test roots, so both test binaries see the same authored descriptors; the two need
# distinct NAMEs.
# distinct NAMEs. The archive_fixture set is also packed twice, once for each binary.
set(_authored "${HIPKERNELPROVIDER_TEST_DESCRIPTOR_SOURCE_ROOT}")
set(_unit "${HIPKERNELPROVIDER_UNIT_BUILD_DIR}")
set(_integration "${HIPKERNELPROVIDER_INTEGRATION_BUILD_DIR}")
Expand Down Expand Up @@ -1711,6 +1711,18 @@ function(hkp_add_packaging)
PACK_JOBS 1
${_root_args})

# The unit binary's copy of the same set. It stages outside the unit root, so the unit
# engine does not discover it, and the integration suite cannot corrupt it.
_hkp_wire_root(
NAME unit_archive_fixture
SOURCE_ROOT "${_authored}/${HIPKERNELPROVIDER_TEST_SET_ARCHIVE_FIXTURE}"
ARCHES "${_arches}"
HIPCC "${HKP_HIPCC}"
ROCM_KPACK_DIR "${_rocm_kpack_dir}"
OUT_ROOT "${HIPKERNELPROVIDER_UNIT_ARCHIVE_FIXTURE_BUILD_DIR}"
PACK_JOBS 1
${_root_args})

hkp_register_tests("${_rocm_kpack_dir}" "${HKP_HIPCC}" "${_rocke_comgr_lib}")
endfunction()

Expand Down
14 changes: 11 additions & 3 deletions dnn-providers/hip-kernel-provider/src/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -114,14 +114,22 @@ if(HIPDNN_ENABLE_KERNEL_INGESTOR)
HIPKERNELPROVIDER_DESCRIPTOR_SUBDIR="${HIPKERNELPROVIDER_DESCRIPTOR_SUBDIR}"
)

# The kpack kernel source adapter. Listed here rather than beside Kernel.cpp above
# because it does not compile without the ingestor: KpackArchive.cpp is the only
# translation unit in the provider that includes <rocm_kpack/kpack.h>.
# The kpack kernel source adapter, and the loader for saved code objects. Listed here
# rather than beside Kernel.cpp above because they do not compile without the
# ingestor: KpackArchive.cpp is the only translation unit in the provider that
# includes <rocm_kpack/kpack.h>.
target_sources(hip_kernel_provider_impl PRIVATE
compilation/CodeObjectLoad.cpp
compilation/KpackArchive.cpp
compilation/KpackKernelLoader.cpp
compilation/KpackProgram.cpp
)

# The plan serialization hooks. They save and load only ingestor plans, so a build
# without the ingestor exports none of them.
target_sources(hip_kernel_provider_impl PRIVATE
core/PluginSerialization.cpp
)
endif()

# hip_kernel_provider - SHARED library (the actual plugin)
Expand Down
Original file line number Diff line number Diff line change
@@ -0,0 +1,76 @@
// Copyright © Advanced Micro Devices, Inc., or its affiliates.
// SPDX-License-Identifier: MIT

#ifdef HIPDNN_ENABLE_KERNEL_INGESTOR

#include "CodeObjectLoad.hpp"

#include "Kernel.hpp"
#include "KpackModule.hpp"
#include "KpackProgram.hpp"
#include "device/ScopedDevice.hpp"

#include <utility>

namespace hip_kernel_provider::compilation
{

LoadedCodeObject loadCodeObjectOnDevice(const std::vector<uint8_t>& bytes,
const std::string& symbol,
int deviceOrdinal,
const std::string& label)
{
// A module belongs to the device that is current at hipModuleLoadData. Make the
// target device current first.
const device::ScopedDevice binding(deviceOrdinal);
if(!binding.bound())
{
throw CodeObjectLoadFailure(CodeObjectLoadStage::BIND_DEVICE,
hipErrorInvalidDevice,
"cannot make device " + std::to_string(deviceOrdinal)
+ " current to load the code object of " + label);
}

// hipModuleLoadData takes no length, so it must not receive an empty buffer.
if(bytes.empty())
{
throw CodeObjectLoadFailure(CodeObjectLoadStage::MODULE_LOAD,
hipErrorInvalidImage,
"hipModuleLoadData was not called for " + label
+ ": the code object is empty");
}

hipModule_t module = nullptr;
hipError_t status = hipModuleLoadData(&module, bytes.data());
if(status != hipSuccess)
{
throw CodeObjectLoadFailure(CodeObjectLoadStage::MODULE_LOAD,
status,
"hipModuleLoadData rejected the code object of " + label
+ " on device " + std::to_string(deviceOrdinal) + ": "
+ hipGetErrorString(status));
}

// Owned before the symbol lookup, so a failed lookup unloads the module on its device.
auto owner = std::make_shared<const KpackModule>(module, deviceOrdinal);

hipFunction_t function = nullptr;
status = hipModuleGetFunction(&function, module, symbol.c_str());
if(status != hipSuccess)
{
throw CodeObjectLoadFailure(CodeObjectLoadStage::SYMBOL_LOOKUP,
status,
"hipModuleGetFunction found no symbol '" + symbol
+ "' in the code object of " + label + ": "
+ hipGetErrorString(status));
}

LoadedCodeObject loaded;
loaded.program = std::make_unique<KpackProgram>(std::move(owner), label);
loaded.kernel = std::make_unique<Kernel>(function, symbol, deviceOrdinal);
return loaded;
}

} // namespace hip_kernel_provider::compilation

#endif // HIPDNN_ENABLE_KERNEL_INGESTOR
Loading
Loading