Factoring out common debug info from GPU executable flatbuffers.
This also adds source file publishing to all GPU targets. Basic support
for export-specific debug info is added but switching targets to use it
is left to a future change.
diff --git a/build_tools/bazel/iree_flatcc.bzl b/build_tools/bazel/iree_flatcc.bzl
index 80e66d6..de5ddb7 100644
--- a/build_tools/bazel/iree_flatcc.bzl
+++ b/build_tools/bazel/iree_flatcc.bzl
@@ -10,12 +10,14 @@
name,
srcs,
flatcc_args = ["--common", "--reader"],
+ includes = [],
testonly = False,
**kwargs):
flatcc = "@com_github_dvidelabs_flatcc//:flatcc"
flags = [
"-o$(RULEDIR)",
+ "-I runtime/src",
] + flatcc_args
out_stem = "%s" % (srcs[0].replace(".fbs", ""))
@@ -34,10 +36,10 @@
native.genrule(
name = name + "_gen",
- srcs = srcs,
+ srcs = srcs + includes,
outs = outs,
tools = [flatcc],
- cmd = "$(location %s) %s $(SRCS)" % (flatcc, " ".join(flags)),
+ cmd = "$(location %s) %s %s" % (flatcc, " ".join(flags), " ".join(["$(location {})".format(src) for src in srcs])),
testonly = testonly,
)
native.cc_library(
diff --git a/build_tools/bazel_to_cmake/bazel_to_cmake_converter.py b/build_tools/bazel_to_cmake/bazel_to_cmake_converter.py
index 9a1796b..eb5d2b1 100644
--- a/build_tools/bazel_to_cmake/bazel_to_cmake_converter.py
+++ b/build_tools/bazel_to_cmake/bazel_to_cmake_converter.py
@@ -662,16 +662,18 @@
f" PUBLIC\n)\n\n"
)
- def iree_flatbuffer_c_library(self, name, srcs, flatcc_args=None):
+ def iree_flatbuffer_c_library(self, name, srcs, flatcc_args=None, includes=None):
name_block = self._convert_string_arg_block("NAME", name, quote=False)
srcs_block = self._convert_srcs_block(srcs)
flatcc_args_block = self._convert_string_list_block("FLATCC_ARGS", flatcc_args)
+ includes_block = self._convert_srcs_block(includes, block_name="INCLUDES")
self._converter.body += (
f"flatbuffer_c_library(\n"
f"{name_block}"
f"{srcs_block}"
f"{flatcc_args_block}"
+ f"{includes_block}"
f" PUBLIC\n)\n\n"
)
diff --git a/build_tools/cmake/flatbuffer_c_library.cmake b/build_tools/cmake/flatbuffer_c_library.cmake
index fe0913c..2016cdf 100644
--- a/build_tools/cmake/flatbuffer_c_library.cmake
+++ b/build_tools/cmake/flatbuffer_c_library.cmake
@@ -48,7 +48,7 @@
cmake_parse_arguments(_RULE
"PUBLIC;TESTONLY"
"NAME"
- "SRCS;FLATCC_ARGS"
+ "SRCS;FLATCC_ARGS;INCLUDES"
${ARGN}
)
@@ -94,6 +94,7 @@
iree-flatcc-cli
-o "${CMAKE_CURRENT_BINARY_DIR}"
-I "${IREE_ROOT_DIR}"
+ -I "${IREE_ROOT_DIR}/runtime/src"
${_RULE_FLATCC_ARGS}
"${_RULE_SRCS}"
WORKING_DIRECTORY
diff --git a/compiler/plugins/target/CUDA/BUILD.bazel b/compiler/plugins/target/CUDA/BUILD.bazel
index 2d475bf..b694187 100644
--- a/compiler/plugins/target/CUDA/BUILD.bazel
+++ b/compiler/plugins/target/CUDA/BUILD.bazel
@@ -33,11 +33,13 @@
"//compiler/src/iree/compiler/Codegen/LLVMGPU",
"//compiler/src/iree/compiler/Codegen/Utils",
"//compiler/src/iree/compiler/Dialect/HAL/Target",
+ "//compiler/src/iree/compiler/Dialect/HAL/Utils:ExecutableDebugInfoUtils",
"//compiler/src/iree/compiler/Dialect/HAL/Utils:LLVMLinkerUtils",
"//compiler/src/iree/compiler/PluginAPI",
"//compiler/src/iree/compiler/Utils",
"//runtime/src/iree/base/internal/flatcc:building",
"//runtime/src/iree/schemas:cuda_executable_def_c_fbs",
+ "//runtime/src/iree/schemas:executable_debug_info_c_fbs",
"@iree_cuda//:libdevice_embedded",
"@llvm-project//llvm:Analysis",
"@llvm-project//llvm:BitReader",
diff --git a/compiler/plugins/target/CUDA/CMakeLists.txt b/compiler/plugins/target/CUDA/CMakeLists.txt
index 214f78b..70c6dc6 100644
--- a/compiler/plugins/target/CUDA/CMakeLists.txt
+++ b/compiler/plugins/target/CUDA/CMakeLists.txt
@@ -57,10 +57,12 @@
iree::compiler::Codegen::LLVMGPU
iree::compiler::Codegen::Utils
iree::compiler::Dialect::HAL::Target
+ iree::compiler::Dialect::HAL::Utils::ExecutableDebugInfoUtils
iree::compiler::Dialect::HAL::Utils::LLVMLinkerUtils
iree::compiler::PluginAPI
iree::compiler::Utils
iree::schemas::cuda_executable_def_c_fbs
+ iree::schemas::executable_debug_info_c_fbs
iree_cuda::libdevice_embedded
PUBLIC
)
diff --git a/compiler/plugins/target/CUDA/CUDATarget.cpp b/compiler/plugins/target/CUDA/CUDATarget.cpp
index 14e6a32..47c785c 100644
--- a/compiler/plugins/target/CUDA/CUDATarget.cpp
+++ b/compiler/plugins/target/CUDA/CUDATarget.cpp
@@ -10,6 +10,7 @@
#include "iree/compiler/Codegen/LLVMGPU/Passes.h"
#include "iree/compiler/Codegen/Utils/GPUUtils.h"
#include "iree/compiler/Dialect/HAL/Target/TargetRegistry.h"
+#include "iree/compiler/Dialect/HAL/Utils/ExecutableDebugInfoUtils.h"
#include "iree/compiler/Dialect/HAL/Utils/LLVMLinkerUtils.h"
#include "iree/compiler/PluginAPI/Client.h"
#include "iree/compiler/Utils/FlatbufferUtils.h"
@@ -522,9 +523,13 @@
FlatbufferBuilder builder;
iree_hal_cuda_ExecutableDef_start_as_root(builder);
+ // Attach embedded source file contents.
+ auto sourceFilesRef = createSourceFilesVec(
+ serOptions.debugLevel, variantOp.getSourcesAttr(), builder);
+
SmallVector<std::string> entryPointNames;
std::string ptxImage;
- SmallVector<iree_hal_cuda_FileLineLocDef_ref_t> sourceLocationRefs;
+ SmallVector<iree_hal_debug_FileLineLocDef_ref_t> sourceLocationRefs;
if (variantOp.isExternal()) {
if (!variantOp.getObjects().has_value()) {
return variantOp.emitOpError()
@@ -595,7 +600,7 @@
if (serOptions.debugLevel >= 1) {
if (auto loc = findFirstFileLoc(exportOp.getLoc())) {
auto filenameRef = builder.createString(loc->getFilename());
- sourceLocationRefs.push_back(iree_hal_cuda_FileLineLocDef_create(
+ sourceLocationRefs.push_back(iree_hal_debug_FileLineLocDef_create(
builder, filenameRef, loc->getLine()));
}
}
@@ -691,6 +696,7 @@
iree_hal_cuda_ExecutableDef_source_locations_add(builder,
sourceLocationsRef);
}
+ iree_hal_cuda_ExecutableDef_source_files_add(builder, sourceFilesRef);
iree_hal_cuda_ExecutableDef_end_as_root(builder);
// Add the binary data to the target executable.
diff --git a/compiler/plugins/target/MetalSPIRV/BUILD.bazel b/compiler/plugins/target/MetalSPIRV/BUILD.bazel
index ede5566..ae750cb 100644
--- a/compiler/plugins/target/MetalSPIRV/BUILD.bazel
+++ b/compiler/plugins/target/MetalSPIRV/BUILD.bazel
@@ -31,8 +31,10 @@
"//compiler/src/iree/compiler/Codegen/Utils",
"//compiler/src/iree/compiler/Dialect/Flow/IR",
"//compiler/src/iree/compiler/Dialect/HAL/Target",
+ "//compiler/src/iree/compiler/Dialect/HAL/Utils:ExecutableDebugInfoUtils",
"//compiler/src/iree/compiler/PluginAPI",
"//compiler/src/iree/compiler/Utils",
+ "//runtime/src/iree/schemas:executable_debug_info_c_fbs",
"//runtime/src/iree/schemas:metal_executable_def_c_fbs",
"@llvm-project//llvm:Support",
"@llvm-project//llvm:TargetParser",
diff --git a/compiler/plugins/target/MetalSPIRV/CMakeLists.txt b/compiler/plugins/target/MetalSPIRV/CMakeLists.txt
index 678a37a..3a7b0e6 100644
--- a/compiler/plugins/target/MetalSPIRV/CMakeLists.txt
+++ b/compiler/plugins/target/MetalSPIRV/CMakeLists.txt
@@ -41,8 +41,10 @@
iree::compiler::Codegen::Utils
iree::compiler::Dialect::Flow::IR
iree::compiler::Dialect::HAL::Target
+ iree::compiler::Dialect::HAL::Utils::ExecutableDebugInfoUtils
iree::compiler::PluginAPI
iree::compiler::Utils
+ iree::schemas::executable_debug_info_c_fbs
iree::schemas::metal_executable_def_c_fbs
PUBLIC
)
diff --git a/compiler/plugins/target/MetalSPIRV/MetalSPIRVTarget.cpp b/compiler/plugins/target/MetalSPIRV/MetalSPIRVTarget.cpp
index 4fa2b03..8372307 100644
--- a/compiler/plugins/target/MetalSPIRV/MetalSPIRVTarget.cpp
+++ b/compiler/plugins/target/MetalSPIRV/MetalSPIRVTarget.cpp
@@ -12,6 +12,7 @@
#include "iree/compiler/Codegen/SPIRV/Passes.h"
#include "iree/compiler/Dialect/Flow/IR/FlowDialect.h"
#include "iree/compiler/Dialect/HAL/Target/TargetRegistry.h"
+#include "iree/compiler/Dialect/HAL/Utils/ExecutableDebugInfoUtils.h"
#include "iree/compiler/PluginAPI/Client.h"
#include "iree/compiler/Utils/FlatbufferUtils.h"
#include "iree/schemas/metal_executable_def_builder.h"
@@ -212,6 +213,10 @@
FlatbufferBuilder builder;
iree_hal_metal_ExecutableDef_start_as_root(builder);
+ // Attach embedded source file contents.
+ auto sourceFilesRef = createSourceFilesVec(
+ serOptions.debugLevel, variantOp.getSourcesAttr(), builder);
+
auto entryPointNamesRef = builder.createStringVec(mslEntryPointNames);
iree_hal_metal_ExecutableDef_entry_points_add(builder, entryPointNamesRef);
@@ -243,6 +248,8 @@
iree_hal_metal_ExecutableDef_shader_libraries_add(builder, libsRef);
}
+ iree_hal_metal_ExecutableDef_source_files_add(builder, sourceFilesRef);
+
iree_hal_metal_ExecutableDef_end_as_root(builder);
// 5. Add the binary data to the target executable.
diff --git a/compiler/plugins/target/ROCM/BUILD.bazel b/compiler/plugins/target/ROCM/BUILD.bazel
index 75296b8..4771120 100644
--- a/compiler/plugins/target/ROCM/BUILD.bazel
+++ b/compiler/plugins/target/ROCM/BUILD.bazel
@@ -35,9 +35,11 @@
"//compiler/src/iree/compiler/Codegen/Utils",
"//compiler/src/iree/compiler/Dialect/HAL/IR",
"//compiler/src/iree/compiler/Dialect/HAL/Target",
+ "//compiler/src/iree/compiler/Dialect/HAL/Utils:ExecutableDebugInfoUtils",
"//compiler/src/iree/compiler/Dialect/HAL/Utils:LLVMLinkerUtils",
"//compiler/src/iree/compiler/PluginAPI",
"//compiler/src/iree/compiler/Utils",
+ "//runtime/src/iree/schemas:executable_debug_info_c_fbs",
"//runtime/src/iree/schemas:rocm_executable_def_c_fbs",
"@llvm-project//llvm:AMDGPUCodeGen",
"@llvm-project//llvm:Analysis",
diff --git a/compiler/plugins/target/ROCM/CMakeLists.txt b/compiler/plugins/target/ROCM/CMakeLists.txt
index b3e8fd5..ca749fe 100644
--- a/compiler/plugins/target/ROCM/CMakeLists.txt
+++ b/compiler/plugins/target/ROCM/CMakeLists.txt
@@ -60,9 +60,11 @@
iree::compiler::Codegen::Utils
iree::compiler::Dialect::HAL::IR
iree::compiler::Dialect::HAL::Target
+ iree::compiler::Dialect::HAL::Utils::ExecutableDebugInfoUtils
iree::compiler::Dialect::HAL::Utils::LLVMLinkerUtils
iree::compiler::PluginAPI
iree::compiler::Utils
+ iree::schemas::executable_debug_info_c_fbs
iree::schemas::rocm_executable_def_c_fbs
PUBLIC
)
diff --git a/compiler/plugins/target/ROCM/ROCMTarget.cpp b/compiler/plugins/target/ROCM/ROCMTarget.cpp
index 62b020c..b601deb 100644
--- a/compiler/plugins/target/ROCM/ROCMTarget.cpp
+++ b/compiler/plugins/target/ROCM/ROCMTarget.cpp
@@ -18,6 +18,7 @@
#include "iree/compiler/Codegen/Utils/Utils.h"
#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
#include "iree/compiler/Dialect/HAL/Target/TargetRegistry.h"
+#include "iree/compiler/Dialect/HAL/Utils/ExecutableDebugInfoUtils.h"
#include "iree/compiler/Dialect/HAL/Utils/LLVMLinkerUtils.h"
#include "iree/compiler/PluginAPI/Client.h"
#include "iree/compiler/Utils/FlatbufferUtils.h"
@@ -607,26 +608,11 @@
iree_hal_rocm_ExecutableDef_start_as_root(builder);
// Attach embedded source file contents.
- SmallVector<iree_hal_rocm_SourceFileDef_ref_t> sourceFileRefs;
- if (auto sourcesAttr = variantOp.getSourcesAttr()) {
- for (auto sourceAttr : llvm::reverse(sourcesAttr.getValue())) {
- if (auto resourceAttr = dyn_cast_if_present<DenseResourceElementsAttr>(
- sourceAttr.getValue())) {
- auto filenameRef = builder.createString(sourceAttr.getName());
- auto contentRef = builder.streamUint8Vec([&](llvm::raw_ostream &os) {
- auto blobData = resourceAttr.getRawHandle().getBlob()->getData();
- os.write(blobData.data(), blobData.size());
- return true;
- });
- sourceFileRefs.push_back(iree_hal_rocm_SourceFileDef_create(
- builder, filenameRef, contentRef));
- }
- }
- std::reverse(sourceFileRefs.begin(), sourceFileRefs.end());
- }
+ auto sourceFilesRef = createSourceFilesVec(
+ serOptions.debugLevel, variantOp.getSourcesAttr(), builder);
SmallVector<StringRef> entryPointNames;
- SmallVector<iree_hal_rocm_FileLineLocDef_ref_t> sourceLocationRefs;
+ SmallVector<iree_hal_debug_FileLineLocDef_ref_t> sourceLocationRefs;
entryPointNames.resize(exportOps.size());
for (auto exportOp : exportOps) {
auto ordinalAttr = exportOp.getOrdinalAttr();
@@ -644,27 +630,28 @@
// be kept as-is.
sourceLocationRefs.resize(exportOps.size());
auto filenameRef = builder.createString(loc->getFilename());
- sourceLocationRefs[ordinal] = iree_hal_rocm_FileLineLocDef_create(
+ sourceLocationRefs[ordinal] = iree_hal_debug_FileLineLocDef_create(
builder, filenameRef, loc->getLine());
}
}
}
// Optional compilation stage source files.
- SmallVector<iree_hal_rocm_StageLocationsDef_ref_t> stageLocationsRefs;
+ SmallVector<iree_hal_debug_StageLocationsDef_ref_t> stageLocationsRefs;
if (serOptions.debugLevel >= 3) {
for (auto exportOp : exportOps) {
- SmallVector<iree_hal_rocm_StageLocationDef_ref_t> stageLocationRefs;
+ SmallVector<iree_hal_debug_StageLocationDef_ref_t> stageLocationRefs;
if (auto locsAttr = exportOp.getSourceLocsAttr()) {
for (auto locAttr : locsAttr.getValue()) {
if (auto loc =
findFirstFileLoc(cast<LocationAttr>(locAttr.getValue()))) {
auto stageNameRef = builder.createString(locAttr.getName());
auto filenameRef = builder.createString(loc->getFilename());
- stageLocationRefs.push_back(iree_hal_rocm_StageLocationDef_create(
- builder, stageNameRef,
- iree_hal_rocm_FileLineLocDef_create(builder, filenameRef,
- loc->getLine())));
+ stageLocationRefs.push_back(
+ iree_hal_debug_StageLocationDef_create(
+ builder, stageNameRef,
+ iree_hal_debug_FileLineLocDef_create(builder, filenameRef,
+ loc->getLine())));
}
}
}
@@ -673,7 +660,7 @@
// be kept as-is.
stageLocationsRefs.resize(exportOps.size());
int64_t ordinal = exportOp.getOrdinalAttr().getInt();
- stageLocationsRefs[ordinal] = iree_hal_rocm_StageLocationsDef_create(
+ stageLocationsRefs[ordinal] = iree_hal_debug_StageLocationsDef_create(
builder, builder.createOffsetVecDestructive(stageLocationRefs));
}
}
@@ -710,10 +697,7 @@
iree_hal_rocm_ExecutableDef_stage_locations_add(builder,
stageLocationsRef);
}
- if (!sourceFileRefs.empty()) {
- auto sourceFilesRef = builder.createOffsetVecDestructive(sourceFileRefs);
- iree_hal_rocm_ExecutableDef_source_files_add(builder, sourceFilesRef);
- }
+ iree_hal_rocm_ExecutableDef_source_files_add(builder, sourceFilesRef);
iree_hal_rocm_ExecutableDef_end_as_root(builder);
// Add the binary data to the target executable.
diff --git a/compiler/plugins/target/VulkanSPIRV/BUILD.bazel b/compiler/plugins/target/VulkanSPIRV/BUILD.bazel
index 984bef9..8419ed7 100644
--- a/compiler/plugins/target/VulkanSPIRV/BUILD.bazel
+++ b/compiler/plugins/target/VulkanSPIRV/BUILD.bazel
@@ -29,8 +29,10 @@
"//compiler/src/iree/compiler/Codegen/SPIRV",
"//compiler/src/iree/compiler/Codegen/Utils",
"//compiler/src/iree/compiler/Dialect/HAL/Target",
+ "//compiler/src/iree/compiler/Dialect/HAL/Utils:ExecutableDebugInfoUtils",
"//compiler/src/iree/compiler/PluginAPI",
"//compiler/src/iree/compiler/Utils",
+ "//runtime/src/iree/schemas:executable_debug_info_c_fbs",
"//runtime/src/iree/schemas:spirv_executable_def_c_fbs",
"@llvm-project//llvm:Support",
"@llvm-project//mlir:AsmParser",
diff --git a/compiler/plugins/target/VulkanSPIRV/CMakeLists.txt b/compiler/plugins/target/VulkanSPIRV/CMakeLists.txt
index 958e277..b554617 100644
--- a/compiler/plugins/target/VulkanSPIRV/CMakeLists.txt
+++ b/compiler/plugins/target/VulkanSPIRV/CMakeLists.txt
@@ -37,8 +37,10 @@
iree::compiler::Codegen::SPIRV
iree::compiler::Codegen::Utils
iree::compiler::Dialect::HAL::Target
+ iree::compiler::Dialect::HAL::Utils::ExecutableDebugInfoUtils
iree::compiler::PluginAPI
iree::compiler::Utils
+ iree::schemas::executable_debug_info_c_fbs
iree::schemas::spirv_executable_def_c_fbs
PUBLIC
)
diff --git a/compiler/plugins/target/VulkanSPIRV/VulkanSPIRVTarget.cpp b/compiler/plugins/target/VulkanSPIRV/VulkanSPIRVTarget.cpp
index 45bbdf3..10c78e3 100644
--- a/compiler/plugins/target/VulkanSPIRV/VulkanSPIRVTarget.cpp
+++ b/compiler/plugins/target/VulkanSPIRV/VulkanSPIRVTarget.cpp
@@ -8,6 +8,7 @@
#include "iree/compiler/Codegen/Dialect/GPU/TargetUtils/KnownTargets.h"
#include "iree/compiler/Codegen/SPIRV/Passes.h"
#include "iree/compiler/Dialect/HAL/Target/TargetRegistry.h"
+#include "iree/compiler/Dialect/HAL/Utils/ExecutableDebugInfoUtils.h"
#include "iree/compiler/PluginAPI/Client.h"
#include "iree/compiler/Utils/FlatbufferUtils.h"
#include "iree/compiler/Utils/ModuleUtils.h"
@@ -184,23 +185,8 @@
iree_hal_spirv_ExecutableDef_start_as_root(builder);
// Attach embedded source file contents.
- SmallVector<iree_hal_spirv_SourceFileDef_ref_t> sourceFileRefs;
- if (auto sourcesAttr = variantOp.getSourcesAttr()) {
- for (auto sourceAttr : llvm::reverse(sourcesAttr.getValue())) {
- if (auto resourceAttr = dyn_cast_if_present<DenseResourceElementsAttr>(
- sourceAttr.getValue())) {
- auto filenameRef = builder.createString(sourceAttr.getName());
- auto contentRef = builder.streamUint8Vec([&](llvm::raw_ostream &os) {
- auto blobData = resourceAttr.getRawHandle().getBlob()->getData();
- os.write(blobData.data(), blobData.size());
- return true;
- });
- sourceFileRefs.push_back(iree_hal_spirv_SourceFileDef_create(
- builder, filenameRef, contentRef));
- }
- }
- std::reverse(sourceFileRefs.begin(), sourceFileRefs.end());
- }
+ auto sourceFilesRef = createSourceFilesVec(
+ options.debugLevel, variantOp.getSourcesAttr(), builder);
// The list of shader modules.
SmallVector<iree_hal_spirv_ShaderModuleDef_ref_t> shaderModuleRefs;
@@ -211,7 +197,7 @@
SmallVector<StringRef> entryPointNames;
SmallVector<uint32_t> subgroupSizes;
SmallVector<uint32_t> shaderModuleIndices;
- SmallVector<iree_hal_spirv_FileLineLocDef_ref_t> sourceLocationRefs;
+ SmallVector<iree_hal_debug_FileLineLocDef_ref_t> sourceLocationRefs;
entryPointNames.resize(ordinalCount);
subgroupSizes.resize(ordinalCount);
shaderModuleIndices.resize(ordinalCount);
@@ -275,17 +261,17 @@
// kept as-is.
sourceLocationRefs.resize(ordinalCount);
auto filenameRef = builder.createString(loc->getFilename());
- sourceLocationRefs[ordinal] = iree_hal_spirv_FileLineLocDef_create(
+ sourceLocationRefs[ordinal] = iree_hal_debug_FileLineLocDef_create(
builder, filenameRef, loc->getLine());
}
}
}
// Optional compilation stage source files.
- SmallVector<iree_hal_spirv_StageLocationsDef_ref_t> stageLocationsRefs;
+ SmallVector<iree_hal_debug_StageLocationsDef_ref_t> stageLocationsRefs;
if (options.debugLevel >= 3) {
for (auto exportOp : exportOps) {
- SmallVector<iree_hal_spirv_StageLocationDef_ref_t> stageLocationRefs;
+ SmallVector<iree_hal_debug_StageLocationDef_ref_t> stageLocationRefs;
if (auto locsAttr = exportOp.getSourceLocsAttr()) {
for (auto locAttr : locsAttr.getValue()) {
if (auto loc =
@@ -293,9 +279,9 @@
auto stageNameRef = builder.createString(locAttr.getName());
auto filenameRef = builder.createString(loc->getFilename());
stageLocationRefs.push_back(
- iree_hal_spirv_StageLocationDef_create(
+ iree_hal_debug_StageLocationDef_create(
builder, stageNameRef,
- iree_hal_spirv_FileLineLocDef_create(builder, filenameRef,
+ iree_hal_debug_FileLineLocDef_create(builder, filenameRef,
loc->getLine())));
}
}
@@ -305,7 +291,7 @@
// be kept as-is.
stageLocationsRefs.resize(ordinalCount);
int64_t ordinal = exportOp.getOrdinalAttr().getInt();
- stageLocationsRefs[ordinal] = iree_hal_spirv_StageLocationsDef_create(
+ stageLocationsRefs[ordinal] = iree_hal_debug_StageLocationsDef_create(
builder, builder.createOffsetVecDestructive(stageLocationRefs));
}
}
@@ -339,10 +325,7 @@
iree_hal_spirv_ExecutableDef_stage_locations_add(builder,
stageLocationsRef);
}
- if (!sourceFileRefs.empty()) {
- auto sourceFilesRef = builder.createOffsetVecDestructive(sourceFileRefs);
- iree_hal_spirv_ExecutableDef_source_files_add(builder, sourceFilesRef);
- }
+ iree_hal_spirv_ExecutableDef_source_files_add(builder, sourceFilesRef);
iree_hal_spirv_ExecutableDef_end_as_root(builder);
diff --git a/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt b/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt
index caf4460..4b7ef4e 100644
--- a/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt
+++ b/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt
@@ -52,8 +52,10 @@
iree::compiler::Codegen::SPIRV
iree::compiler::Dialect::Flow::IR
iree::compiler::Dialect::HAL::Target
+ iree::compiler::Dialect::HAL::Utils::ExecutableDebugInfoUtils
iree::compiler::PluginAPI
iree::compiler::Utils
+ iree::schemas::executable_debug_info_c_fbs
iree::schemas::wgsl_executable_def_c_fbs
libtint
PUBLIC
diff --git a/compiler/plugins/target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp b/compiler/plugins/target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp
index 8fd7c53..61d9965 100644
--- a/compiler/plugins/target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp
+++ b/compiler/plugins/target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp
@@ -11,6 +11,7 @@
#include "iree/compiler/Codegen/WGSL/Passes.h"
#include "iree/compiler/Dialect/Flow/IR/FlowDialect.h"
#include "iree/compiler/Dialect/HAL/Target/TargetRegistry.h"
+#include "iree/compiler/Dialect/HAL/Utils/ExecutableDebugInfoUtils.h"
#include "iree/compiler/PluginAPI/Client.h"
#include "iree/compiler/Utils/FlatbufferUtils.h"
#include "iree/schemas/wgsl_executable_def_builder.h"
@@ -238,6 +239,10 @@
FlatbufferBuilder builder;
iree_hal_wgsl_ExecutableDef_start_as_root(builder);
+ // Attach embedded source file contents.
+ auto sourceFilesRef = createSourceFilesVec(
+ serOptions.debugLevel, variantOp.getSourcesAttr(), builder);
+
iree_hal_wgsl_ShaderModuleDef_start(builder);
auto wgslRef = builder.createString(wgsl.value());
iree_hal_wgsl_ShaderModuleDef_code_add(builder, wgslRef);
@@ -251,6 +256,7 @@
auto entryPointsRef = flatbuffers_uint32_vec_create(
builder, entryPointOrdinals.data(), entryPointOrdinals.size());
iree_hal_wgsl_ExecutableDef_entry_points_add(builder, entryPointsRef);
+ iree_hal_wgsl_ExecutableDef_source_files_add(builder, sourceFilesRef);
iree_hal_wgsl_ExecutableDef_end_as_root(builder);
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Utils/BUILD.bazel b/compiler/src/iree/compiler/Dialect/HAL/Utils/BUILD.bazel
index 7745815..2a77d86 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Utils/BUILD.bazel
+++ b/compiler/src/iree/compiler/Dialect/HAL/Utils/BUILD.bazel
@@ -13,6 +13,22 @@
)
iree_compiler_cc_library(
+ name = "ExecutableDebugInfoUtils",
+ srcs = [
+ "ExecutableDebugInfoUtils.cpp",
+ ],
+ hdrs = [
+ "ExecutableDebugInfoUtils.h",
+ ],
+ deps = [
+ "//compiler/src/iree/compiler/Dialect/HAL/IR",
+ "//compiler/src/iree/compiler/Utils",
+ "//runtime/src/iree/schemas:executable_debug_info_c_fbs",
+ "@llvm-project//mlir:IR",
+ ],
+)
+
+iree_compiler_cc_library(
name = "LLVMLinkerUtils",
srcs = [
"LLVMLinkerUtils.cpp",
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Utils/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/HAL/Utils/CMakeLists.txt
index 696c5f3..22e7732 100644
--- a/compiler/src/iree/compiler/Dialect/HAL/Utils/CMakeLists.txt
+++ b/compiler/src/iree/compiler/Dialect/HAL/Utils/CMakeLists.txt
@@ -12,6 +12,21 @@
iree_cc_library(
NAME
+ ExecutableDebugInfoUtils
+ HDRS
+ "ExecutableDebugInfoUtils.h"
+ SRCS
+ "ExecutableDebugInfoUtils.cpp"
+ DEPS
+ MLIRIR
+ iree::compiler::Dialect::HAL::IR
+ iree::compiler::Utils
+ iree::schemas::executable_debug_info_c_fbs
+ PUBLIC
+)
+
+iree_cc_library(
+ NAME
LLVMLinkerUtils
HDRS
"LLVMLinkerUtils.h"
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Utils/ExecutableDebugInfoUtils.cpp b/compiler/src/iree/compiler/Dialect/HAL/Utils/ExecutableDebugInfoUtils.cpp
new file mode 100644
index 0000000..cef943e
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/HAL/Utils/ExecutableDebugInfoUtils.cpp
@@ -0,0 +1,99 @@
+// Copyright 2024 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#include "iree/compiler/Dialect/HAL/Utils/ExecutableDebugInfoUtils.h"
+
+#include "iree/compiler/Utils/ModuleUtils.h"
+#include "iree/schemas/executable_debug_info_builder.h"
+#include "mlir/IR/DialectResourceBlobManager.h"
+
+namespace mlir::iree_compiler::IREE::HAL {
+
+flatbuffers_vec_ref_t createSourceFilesVec(int debugLevel,
+ DictionaryAttr sourcesAttr,
+ FlatbufferBuilder &fbb) {
+ if (debugLevel < 1) {
+ // No debug information.
+ return 0;
+ } else if (!sourcesAttr || sourcesAttr.empty()) {
+ // No sources embedded in the IR.
+ return 0;
+ }
+ SmallVector<iree_hal_debug_SourceFileDef_ref_t> sourceFileRefs;
+ for (auto sourceAttr : llvm::reverse(sourcesAttr.getValue())) {
+ if (auto resourceAttr = dyn_cast_if_present<DenseResourceElementsAttr>(
+ sourceAttr.getValue())) {
+ auto filenameRef = fbb.createString(sourceAttr.getName());
+ auto contentRef = fbb.streamUint8Vec([&](llvm::raw_ostream &os) {
+ auto blobData = resourceAttr.getRawHandle().getBlob()->getData();
+ os.write(blobData.data(), blobData.size());
+ return true;
+ });
+ sourceFileRefs.push_back(
+ iree_hal_debug_SourceFileDef_create(fbb, filenameRef, contentRef));
+ }
+ }
+ std::reverse(sourceFileRefs.begin(), sourceFileRefs.end());
+ return fbb.createOffsetVecDestructive(sourceFileRefs);
+}
+
+SmallVector<flatbuffers_ref_t>
+createExportDefs(int debugLevel,
+ ArrayRef<IREE::HAL::ExecutableExportOp> exportOps,
+ FlatbufferBuilder &fbb) {
+ if (debugLevel < 1) {
+ // No debug information.
+ return {};
+ }
+
+ SmallVector<flatbuffers_ref_t> exportDefs;
+ exportDefs.resize(exportOps.size(), 0);
+
+ for (auto exportOp : exportOps) {
+ auto ordinalAttr = exportOp.getOrdinalAttr();
+ assert(ordinalAttr && "ordinals must be assigned");
+ int64_t ordinal = ordinalAttr.getInt();
+
+ flatbuffers_ref_t locationRef = 0;
+ if (debugLevel >= 1) {
+ if (auto loc = findFirstFileLoc(exportOp.getLoc())) {
+ auto filenameRef = fbb.createString(loc->getFilename());
+ locationRef = iree_hal_debug_FileLineLocDef_create(fbb, filenameRef,
+ loc->getLine());
+ }
+ }
+
+ flatbuffers_vec_ref_t stageLocationsRef = 0;
+ if (debugLevel >= 3) {
+ SmallVector<iree_hal_debug_StageLocationDef_ref_t> stageLocationRefs;
+ if (auto locsAttr = exportOp.getSourceLocsAttr()) {
+ for (auto locAttr : locsAttr.getValue()) {
+ if (auto loc =
+ findFirstFileLoc(cast<LocationAttr>(locAttr.getValue()))) {
+ auto stageNameRef = fbb.createString(locAttr.getName());
+ auto filenameRef = fbb.createString(loc->getFilename());
+ stageLocationRefs.push_back(iree_hal_debug_StageLocationDef_create(
+ fbb, stageNameRef,
+ iree_hal_debug_FileLineLocDef_create(fbb, filenameRef,
+ loc->getLine())));
+ }
+ }
+ }
+ if (!stageLocationRefs.empty()) {
+ stageLocationsRef = fbb.createOffsetVecDestructive(stageLocationRefs);
+ }
+ }
+
+ if (locationRef || stageLocationsRef) {
+ exportDefs[ordinal] =
+ iree_hal_debug_ExportDef_create(fbb, locationRef, stageLocationsRef);
+ }
+ }
+
+ return exportDefs;
+}
+
+} // namespace mlir::iree_compiler::IREE::HAL
diff --git a/compiler/src/iree/compiler/Dialect/HAL/Utils/ExecutableDebugInfoUtils.h b/compiler/src/iree/compiler/Dialect/HAL/Utils/ExecutableDebugInfoUtils.h
new file mode 100644
index 0000000..0a6cd02
--- /dev/null
+++ b/compiler/src/iree/compiler/Dialect/HAL/Utils/ExecutableDebugInfoUtils.h
@@ -0,0 +1,43 @@
+// Copyright 2024 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#ifndef IREE_COMPILER_DIALECT_HAL_UTILS_EXECUTABLEDEBUGINFOUTILS_H_
+#define IREE_COMPILER_DIALECT_HAL_UTILS_EXECUTABLEDEBUGINFOUTILS_H_
+
+#include "iree/compiler/Dialect/HAL/IR/HALOps.h"
+#include "iree/compiler/Dialect/HAL/IR/HALTypes.h"
+#include "iree/compiler/Utils/FlatbufferUtils.h"
+
+namespace mlir::iree_compiler::IREE::HAL {
+
+// Creates a `[iree.hal.debug.SourceFileDef]` vector from the given sources
+// dictionary (filename keys to resource elements contents).
+//
+// |debugLevel| generally corresponds to the gcc-style levels 0-3:
+// 0: no debug information
+// 1: minimal debug information
+// 2: default debug information
+// 3: maximal debug information
+flatbuffers_vec_ref_t createSourceFilesVec(int debugLevel,
+ DictionaryAttr sourcesAttr,
+ FlatbufferBuilder &fbb);
+
+// Creates one `iree.hal.debug.ExportDef` for every export and returns them in
+// the same order.
+//
+// |debugLevel| generally corresponds to the gcc-style levels 0-3:
+// 0: no debug information
+// 1: minimal debug information
+// 2: default debug information
+// 3: maximal debug information
+SmallVector<flatbuffers_ref_t>
+createExportDefs(int debugLevel,
+ ArrayRef<IREE::HAL::ExecutableExportOp> exportOps,
+ FlatbufferBuilder &fbb);
+
+} // namespace mlir::iree_compiler::IREE::HAL
+
+#endif // IREE_COMPILER_DIALECT_HAL_UTILS_EXECUTABLEDEBUGINFOUTILS_H_
diff --git a/experimental/webgpu/BUILD.bazel b/experimental/webgpu/BUILD.bazel
index 4e802e6..2906580 100644
--- a/experimental/webgpu/BUILD.bazel
+++ b/experimental/webgpu/BUILD.bazel
@@ -53,8 +53,10 @@
"//runtime/src/iree/hal/drivers/webgpu/platform",
"//runtime/src/iree/hal/drivers/webgpu/shaders",
"//runtime/src/iree/hal/utils:buffer_transfer",
+ "//runtime/src/iree/hal/utils:executable_debug_info",
"//runtime/src/iree/hal/utils:file_transfer",
"//runtime/src/iree/hal/utils:memory_file",
+ "//runtime/src/iree/schemas:executable_debug_info_c_fbs",
"//runtime/src/iree/schemas:wgsl_executable_def_c_fbs",
"@webgpu_headers",
],
diff --git a/experimental/webgpu/executable.c b/experimental/webgpu/executable.c
index ff38225..1929e41 100644
--- a/experimental/webgpu/executable.c
+++ b/experimental/webgpu/executable.c
@@ -10,9 +10,12 @@
#include "iree/base/api.h"
#include "iree/base/internal/inline_array.h"
+#include "iree/hal/utils/executable_debug_info.h"
// flatcc schemas:
#include "iree/base/internal/flatcc/parsing.h"
+#include "iree/schemas/executable_debug_info_reader.h"
+#include "iree/schemas/executable_debug_info_verifier.h"
#include "iree/schemas/wgsl_executable_def_reader.h"
#include "iree/schemas/wgsl_executable_def_verifier.h"
@@ -268,6 +271,10 @@
executable->host_allocator = host_allocator;
executable->entry_point_count = executable_params->pipeline_layout_count;
+ // Publish any embedded source files to the tracing infrastructure.
+ iree_hal_debug_publish_source_files(
+ iree_hal_rocm_ExecutableDef_source_files_get(executable_def));
+
// Create one pipeline per entry point.
flatbuffers_uint32_vec_t entry_points_vec =
iree_hal_wgsl_ExecutableDef_entry_points_get(executable_def);
diff --git a/runtime/src/iree/hal/drivers/cuda/BUILD.bazel b/runtime/src/iree/hal/drivers/cuda/BUILD.bazel
index c0e7069..f34a5d7 100644
--- a/runtime/src/iree/hal/drivers/cuda/BUILD.bazel
+++ b/runtime/src/iree/hal/drivers/cuda/BUILD.bazel
@@ -63,12 +63,14 @@
"//runtime/src/iree/hal/utils:collective_batch",
"//runtime/src/iree/hal/utils:deferred_command_buffer",
"//runtime/src/iree/hal/utils:deferred_work_queue",
+ "//runtime/src/iree/hal/utils:executable_debug_info",
"//runtime/src/iree/hal/utils:file_transfer",
"//runtime/src/iree/hal/utils:memory_file",
"//runtime/src/iree/hal/utils:resource_set",
"//runtime/src/iree/hal/utils:semaphore_base",
"//runtime/src/iree/hal/utils:stream_tracing",
"//runtime/src/iree/schemas:cuda_executable_def_c_fbs",
+ "//runtime/src/iree/schemas:executable_debug_info_c_fbs",
],
)
diff --git a/runtime/src/iree/hal/drivers/cuda/CMakeLists.txt b/runtime/src/iree/hal/drivers/cuda/CMakeLists.txt
index e5f4c67..e2fc8cb 100644
--- a/runtime/src/iree/hal/drivers/cuda/CMakeLists.txt
+++ b/runtime/src/iree/hal/drivers/cuda/CMakeLists.txt
@@ -60,12 +60,14 @@
iree::hal::utils::collective_batch
iree::hal::utils::deferred_command_buffer
iree::hal::utils::deferred_work_queue
+ iree::hal::utils::executable_debug_info
iree::hal::utils::file_transfer
iree::hal::utils::memory_file
iree::hal::utils::resource_set
iree::hal::utils::semaphore_base
iree::hal::utils::stream_tracing
iree::schemas::cuda_executable_def_c_fbs
+ iree::schemas::executable_debug_info_c_fbs
PUBLIC
)
diff --git a/runtime/src/iree/hal/drivers/cuda/native_executable.c b/runtime/src/iree/hal/drivers/cuda/native_executable.c
index 18ab1e6..c3b32e0 100644
--- a/runtime/src/iree/hal/drivers/cuda/native_executable.c
+++ b/runtime/src/iree/hal/drivers/cuda/native_executable.c
@@ -12,11 +12,14 @@
#include "iree/hal/drivers/cuda/cuda_dynamic_symbols.h"
#include "iree/hal/drivers/cuda/cuda_status_util.h"
#include "iree/hal/drivers/cuda/pipeline_layout.h"
+#include "iree/hal/utils/executable_debug_info.h"
// flatcc schemas:
#include "iree/base/internal/flatcc/parsing.h"
#include "iree/schemas/cuda_executable_def_reader.h"
#include "iree/schemas/cuda_executable_def_verifier.h"
+#include "iree/schemas/executable_debug_info_reader.h"
+#include "iree/schemas/executable_debug_info_verifier.h"
typedef struct iree_hal_cuda_native_executable_t {
// Abstract resource used for injecting reference counting and vtable;
@@ -192,6 +195,13 @@
executable->symbols = symbols;
executable->cu_module = module;
executable->entry_point_count = entry_point_count;
+
+ // Publish any embedded source files to the tracing infrastructure.
+ if (iree_status_is_ok(status)) {
+ iree_hal_debug_publish_source_files(
+ iree_hal_cuda_ExecutableDef_source_files_get(executable_def));
+ }
+
for (iree_host_size_t i = 0; i < entry_point_count; i++) {
// Lookup the function in the module; this should always succeed but we
// cannot trust that the input was generated by our compiler.
@@ -263,13 +273,13 @@
IREE_TRACE({
if (iree_hal_cuda_ExecutableDef_source_locations_is_present(
executable_def)) {
- iree_hal_cuda_FileLineLocDef_vec_t source_locs_vec =
+ iree_hal_debug_FileLineLocDef_vec_t source_locs_vec =
iree_hal_cuda_ExecutableDef_source_locations_get(executable_def);
- iree_hal_cuda_FileLineLocDef_table_t source_loc =
- iree_hal_cuda_FileLineLocDef_vec_at(source_locs_vec, i);
+ iree_hal_debug_FileLineLocDef_table_t source_loc =
+ iree_hal_debug_FileLineLocDef_vec_at(source_locs_vec, i);
flatbuffers_string_t filename =
- iree_hal_cuda_FileLineLocDef_filename_get(source_loc);
- uint32_t line = iree_hal_cuda_FileLineLocDef_line_get(source_loc);
+ iree_hal_debug_FileLineLocDef_filename_get(source_loc);
+ uint32_t line = iree_hal_debug_FileLineLocDef_line_get(source_loc);
info->source_filename =
iree_make_string_view(filename, flatbuffers_string_len(filename));
info->source_line = line;
diff --git a/runtime/src/iree/hal/drivers/hip/CMakeLists.txt b/runtime/src/iree/hal/drivers/hip/CMakeLists.txt
index cb9b29e..8490824 100644
--- a/runtime/src/iree/hal/drivers/hip/CMakeLists.txt
+++ b/runtime/src/iree/hal/drivers/hip/CMakeLists.txt
@@ -64,6 +64,7 @@
iree::base::internal::flatcc::parsing
iree::hal
iree::hal::utils::collective_batch
+ iree::hal::utils::executable_debug_info
iree::hal::utils::deferred_command_buffer
iree::hal::utils::deferred_work_queue
iree::hal::utils::file_transfer
@@ -71,6 +72,7 @@
iree::hal::utils::resource_set
iree::hal::utils::semaphore_base
iree::hal::utils::stream_tracing
+ iree::schemas::executable_debug_info_c_fbs
iree::schemas::rocm_executable_def_c_fbs
PUBLIC
)
diff --git a/runtime/src/iree/hal/drivers/hip/native_executable.c b/runtime/src/iree/hal/drivers/hip/native_executable.c
index 5feab82..9e75e12 100644
--- a/runtime/src/iree/hal/drivers/hip/native_executable.c
+++ b/runtime/src/iree/hal/drivers/hip/native_executable.c
@@ -12,10 +12,12 @@
#include "iree/hal/drivers/hip/dynamic_symbols.h"
#include "iree/hal/drivers/hip/pipeline_layout.h"
#include "iree/hal/drivers/hip/status_util.h"
+#include "iree/hal/utils/executable_debug_info.h"
// flatcc schemas:
#include "iree/base/internal/flatcc/parsing.h"
-// Using the existing ROCM schema fow now.
+#include "iree/schemas/executable_debug_info_reader.h"
+#include "iree/schemas/executable_debug_info_verifier.h"
#include "iree/schemas/rocm_executable_def_reader.h"
#include "iree/schemas/rocm_executable_def_verifier.h"
@@ -208,6 +210,11 @@
executable->symbols = symbols;
executable->hip_module = module;
executable->entry_point_count = entry_point_count;
+
+ // Publish any embedded source files to the tracing infrastructure.
+ iree_hal_debug_publish_source_files(
+ iree_hal_rocm_ExecutableDef_source_files_get(executable_def));
+
for (iree_host_size_t i = 0; i < entry_point_count; i++) {
// Lookup the function in the module; this should always succeed but we
// cannot trust that the input was generated by our compiler.
@@ -282,13 +289,13 @@
IREE_TRACE({
if (iree_hal_rocm_ExecutableDef_source_locations_is_present(
executable_def)) {
- iree_hal_rocm_FileLineLocDef_vec_t source_locs_vec =
+ iree_hal_debug_FileLineLocDef_vec_t source_locs_vec =
iree_hal_rocm_ExecutableDef_source_locations_get(executable_def);
- iree_hal_rocm_FileLineLocDef_table_t source_loc =
- iree_hal_rocm_FileLineLocDef_vec_at(source_locs_vec, i);
+ iree_hal_debug_FileLineLocDef_table_t source_loc =
+ iree_hal_debug_FileLineLocDef_vec_at(source_locs_vec, i);
flatbuffers_string_t filename =
- iree_hal_rocm_FileLineLocDef_filename_get(source_loc);
- uint32_t line = iree_hal_rocm_FileLineLocDef_line_get(source_loc);
+ iree_hal_debug_FileLineLocDef_filename_get(source_loc);
+ uint32_t line = iree_hal_debug_FileLineLocDef_line_get(source_loc);
kernel_info->source_filename =
iree_make_string_view(filename, flatbuffers_string_len(filename));
kernel_info->source_line = line;
diff --git a/runtime/src/iree/hal/drivers/metal/CMakeLists.txt b/runtime/src/iree/hal/drivers/metal/CMakeLists.txt
index 85356f3..c4a9835 100644
--- a/runtime/src/iree/hal/drivers/metal/CMakeLists.txt
+++ b/runtime/src/iree/hal/drivers/metal/CMakeLists.txt
@@ -42,9 +42,11 @@
iree::hal
iree::hal::drivers::metal::builtin
iree::hal::utils::deferred_command_buffer
+ iree::hal::utils::executable_debug_info
iree::hal::utils::file_transfer
iree::hal::utils::memory_file
iree::hal::utils::resource_set
+ iree::schemas::executable_debug_info_c_fbs
iree::schemas::metal_executable_def_c_fbs
"-framework Foundation"
"-framework Metal"
diff --git a/runtime/src/iree/hal/drivers/metal/executable.m b/runtime/src/iree/hal/drivers/metal/executable.m
index b30ea15..0416960 100644
--- a/runtime/src/iree/hal/drivers/metal/executable.m
+++ b/runtime/src/iree/hal/drivers/metal/executable.m
@@ -9,9 +9,12 @@
#include <stddef.h>
#include "iree/base/api.h"
+#include "iree/hal/utils/executable_debug_info.h"
// flatcc schemas:
#include "iree/base/internal/flatcc/parsing.h"
+#include "iree/schemas/executable_debug_info_reader.h"
+#include "iree/schemas/executable_debug_info_verifier.h"
#include "iree/schemas/metal_executable_def_reader.h"
#include "iree/schemas/metal_executable_def_verifier.h"
@@ -280,6 +283,12 @@
executable->host_allocator = host_allocator;
executable->entry_point_count = entry_point_count;
+ // Publish any embedded source files to the tracing infrastructure.
+ if (iree_status_is_ok(status)) {
+ iree_hal_debug_publish_source_files(
+ iree_hal_metal_ExecutableDef_source_files_get(executable_def));
+ }
+
size_t shader_library_count = flatbuffers_string_vec_len(shader_libraries_vec);
size_t shader_source_count = flatbuffers_string_vec_len(shader_sources_vec);
diff --git a/runtime/src/iree/hal/drivers/vulkan/BUILD.bazel b/runtime/src/iree/hal/drivers/vulkan/BUILD.bazel
index ce5b68b..68aef34 100644
--- a/runtime/src/iree/hal/drivers/vulkan/BUILD.bazel
+++ b/runtime/src/iree/hal/drivers/vulkan/BUILD.bazel
@@ -79,10 +79,12 @@
"//runtime/src/iree/hal/drivers/vulkan/util:intrusive_list",
"//runtime/src/iree/hal/drivers/vulkan/util:ref_ptr",
"//runtime/src/iree/hal/utils:deferred_command_buffer",
+ "//runtime/src/iree/hal/utils:executable_debug_info",
"//runtime/src/iree/hal/utils:file_transfer",
"//runtime/src/iree/hal/utils:memory_file",
"//runtime/src/iree/hal/utils:resource_set",
"//runtime/src/iree/hal/utils:semaphore_base",
+ "//runtime/src/iree/schemas:executable_debug_info_c_fbs",
"//runtime/src/iree/schemas:spirv_executable_def_c_fbs",
"@vulkan_headers",
],
diff --git a/runtime/src/iree/hal/drivers/vulkan/CMakeLists.txt b/runtime/src/iree/hal/drivers/vulkan/CMakeLists.txt
index 76e376f..b495ae6 100644
--- a/runtime/src/iree/hal/drivers/vulkan/CMakeLists.txt
+++ b/runtime/src/iree/hal/drivers/vulkan/CMakeLists.txt
@@ -74,10 +74,12 @@
iree::hal::drivers::vulkan::util::intrusive_list
iree::hal::drivers::vulkan::util::ref_ptr
iree::hal::utils::deferred_command_buffer
+ iree::hal::utils::executable_debug_info
iree::hal::utils::file_transfer
iree::hal::utils::memory_file
iree::hal::utils::resource_set
iree::hal::utils::semaphore_base
+ iree::schemas::executable_debug_info_c_fbs
iree::schemas::spirv_executable_def_c_fbs
PUBLIC
)
diff --git a/runtime/src/iree/hal/drivers/vulkan/native_executable.cc b/runtime/src/iree/hal/drivers/vulkan/native_executable.cc
index ebfd006..b44efe8 100644
--- a/runtime/src/iree/hal/drivers/vulkan/native_executable.cc
+++ b/runtime/src/iree/hal/drivers/vulkan/native_executable.cc
@@ -16,9 +16,12 @@
#include "iree/hal/drivers/vulkan/native_pipeline_layout.h"
#include "iree/hal/drivers/vulkan/status_util.h"
#include "iree/hal/drivers/vulkan/util/ref_ptr.h"
+#include "iree/hal/utils/executable_debug_info.h"
// flatcc schemas:
#include "iree/base/internal/flatcc/parsing.h"
+#include "iree/schemas/executable_debug_info_reader.h"
+#include "iree/schemas/executable_debug_info_verifier.h"
#include "iree/schemas/spirv_executable_def_reader.h"
#include "iree/schemas/spirv_executable_def_verifier.h"
@@ -30,8 +33,8 @@
iree_string_view_t name;
// Optional debug information.
- IREE_TRACE(iree_hal_spirv_FileLineLocDef_table_t source_location;)
- IREE_TRACE(iree_hal_spirv_StageLocationDef_vec_t stage_locations;)
+ IREE_TRACE(iree_hal_debug_FileLineLocDef_table_t source_location;)
+ IREE_TRACE(iree_hal_debug_StageLocationDef_vec_t stage_locations;)
} iree_hal_vulkan_entry_point_t;
static iree_status_t iree_hal_vulkan_create_shader_module(
@@ -414,44 +417,32 @@
}
}
+ // Publish any embedded source files to the tracing infrastructure.
+ if (iree_status_is_ok(status)) {
+ iree_hal_debug_publish_source_files(
+ iree_hal_spirv_ExecutableDef_source_files_get(executable_def));
+ }
+
#if IREE_TRACING_FEATURES & IREE_TRACING_FEATURE_INSTRUMENTATION
if (iree_status_is_ok(status)) {
if (iree_hal_spirv_ExecutableDef_source_locations_is_present(
executable_def)) {
- iree_hal_spirv_FileLineLocDef_vec_t source_locations_vec =
+ iree_hal_debug_FileLineLocDef_vec_t source_locations_vec =
iree_hal_spirv_ExecutableDef_source_locations_get(executable_def);
for (iree_host_size_t i = 0; i < entry_point_count; ++i) {
executable->entry_points[i].source_location =
- iree_hal_spirv_FileLineLocDef_vec_at(source_locations_vec, i);
+ iree_hal_debug_FileLineLocDef_vec_at(source_locations_vec, i);
}
}
if (iree_hal_spirv_ExecutableDef_stage_locations_is_present(
executable_def)) {
- iree_hal_spirv_StageLocationsDef_vec_t stage_locations_vec =
+ iree_hal_debug_StageLocationsDef_vec_t stage_locations_vec =
iree_hal_spirv_ExecutableDef_stage_locations_get(executable_def);
for (iree_host_size_t i = 0; i < entry_point_count; ++i) {
- iree_hal_spirv_StageLocationsDef_table_t stage_locations =
- iree_hal_spirv_StageLocationsDef_vec_at(stage_locations_vec, i);
+ iree_hal_debug_StageLocationsDef_table_t stage_locations =
+ iree_hal_debug_StageLocationsDef_vec_at(stage_locations_vec, i);
executable->entry_points[i].stage_locations =
- iree_hal_spirv_StageLocationsDef_locations_get(stage_locations);
- }
- }
-
- // Publish any embedded source files to the tracing infrastructure.
- if (iree_hal_spirv_ExecutableDef_source_files_is_present(executable_def)) {
- iree_hal_spirv_SourceFileDef_vec_t source_files_vec =
- iree_hal_spirv_ExecutableDef_source_files_get(executable_def);
- for (iree_host_size_t i = 0;
- i < iree_hal_spirv_SourceFileDef_vec_len(source_files_vec); ++i) {
- iree_hal_spirv_SourceFileDef_table_t source_file =
- iree_hal_spirv_SourceFileDef_vec_at(source_files_vec, i);
- flatbuffers_string_t path =
- iree_hal_spirv_SourceFileDef_path_get(source_file);
- flatbuffers_uint8_vec_t content =
- iree_hal_spirv_SourceFileDef_content_get(source_file);
- IREE_TRACE_PUBLISH_SOURCE_FILE(path, flatbuffers_string_len(path),
- content,
- flatbuffers_uint8_vec_len(content));
+ iree_hal_debug_StageLocationsDef_locations_get(stage_locations);
}
}
}
@@ -500,29 +491,29 @@
out_source_location->func_name = entry_point->name;
#if IREE_TRACING_FEATURES & IREE_TRACING_FEATURE_INSTRUMENTATION
- iree_hal_spirv_FileLineLocDef_table_t source_location =
+ iree_hal_debug_FileLineLocDef_table_t source_location =
entry_point->source_location;
if (entry_point->stage_locations) {
- for (size_t i = 0; i < iree_hal_spirv_StageLocationDef_vec_len(
+ for (size_t i = 0; i < iree_hal_debug_StageLocationDef_vec_len(
entry_point->stage_locations);
++i) {
- iree_hal_spirv_StageLocationDef_table_t stage_location =
- iree_hal_spirv_StageLocationDef_vec_at(entry_point->stage_locations,
+ iree_hal_debug_StageLocationDef_table_t stage_location =
+ iree_hal_debug_StageLocationDef_vec_at(entry_point->stage_locations,
i);
// TODO(benvanik): a way to select what location is chosen. For now we
// just pick the first one.
source_location =
- iree_hal_spirv_StageLocationDef_location_get(stage_location);
+ iree_hal_debug_StageLocationDef_location_get(stage_location);
break;
}
}
if (source_location) {
flatbuffers_string_t filename =
- iree_hal_spirv_FileLineLocDef_filename_get(source_location);
+ iree_hal_debug_FileLineLocDef_filename_get(source_location);
out_source_location->file_name =
iree_make_string_view(filename, flatbuffers_string_len(filename));
out_source_location->line =
- iree_hal_spirv_FileLineLocDef_line_get(source_location);
+ iree_hal_debug_FileLineLocDef_line_get(source_location);
} else {
out_source_location->file_name = out_source_location->func_name;
out_source_location->line = 0;
diff --git a/runtime/src/iree/hal/utils/BUILD.bazel b/runtime/src/iree/hal/utils/BUILD.bazel
index 1531601..0ad6080 100644
--- a/runtime/src/iree/hal/utils/BUILD.bazel
+++ b/runtime/src/iree/hal/utils/BUILD.bazel
@@ -26,16 +26,6 @@
)
iree_runtime_cc_library(
- name = "debug_allocator",
- srcs = ["debug_allocator.c"],
- hdrs = ["debug_allocator.h"],
- deps = [
- "//runtime/src/iree/base",
- "//runtime/src/iree/hal",
- ],
-)
-
-iree_runtime_cc_library(
name = "collective_batch",
srcs = ["collective_batch.c"],
hdrs = ["collective_batch.h"],
@@ -59,6 +49,16 @@
)
iree_runtime_cc_library(
+ name = "debug_allocator",
+ srcs = ["debug_allocator.c"],
+ hdrs = ["debug_allocator.h"],
+ deps = [
+ "//runtime/src/iree/base",
+ "//runtime/src/iree/hal",
+ ],
+)
+
+iree_runtime_cc_library(
name = "deferred_command_buffer",
srcs = ["deferred_command_buffer.c"],
hdrs = ["deferred_command_buffer.h"],
@@ -71,6 +71,17 @@
)
iree_runtime_cc_library(
+ name = "executable_debug_info",
+ srcs = ["executable_debug_info.c"],
+ hdrs = ["executable_debug_info.h"],
+ deps = [
+ "//runtime/src/iree/base",
+ "//runtime/src/iree/base/internal/flatcc:parsing",
+ "//runtime/src/iree/schemas:executable_debug_info_c_fbs",
+ ],
+)
+
+iree_runtime_cc_library(
name = "file_cache",
srcs = ["file_cache.c"],
hdrs = ["file_cache.h"],
diff --git a/runtime/src/iree/hal/utils/CMakeLists.txt b/runtime/src/iree/hal/utils/CMakeLists.txt
index fed7908..4bb9c8c 100644
--- a/runtime/src/iree/hal/utils/CMakeLists.txt
+++ b/runtime/src/iree/hal/utils/CMakeLists.txt
@@ -27,19 +27,6 @@
iree_cc_library(
NAME
- debug_allocator
- HDRS
- "debug_allocator.h"
- SRCS
- "debug_allocator.c"
- DEPS
- iree::base
- iree::hal
- PUBLIC
-)
-
-iree_cc_library(
- NAME
collective_batch
HDRS
"collective_batch.h"
@@ -69,6 +56,19 @@
iree_cc_library(
NAME
+ debug_allocator
+ HDRS
+ "debug_allocator.h"
+ SRCS
+ "debug_allocator.c"
+ DEPS
+ iree::base
+ iree::hal
+ PUBLIC
+)
+
+iree_cc_library(
+ NAME
deferred_command_buffer
HDRS
"deferred_command_buffer.h"
@@ -84,6 +84,20 @@
iree_cc_library(
NAME
+ executable_debug_info
+ HDRS
+ "executable_debug_info.h"
+ SRCS
+ "executable_debug_info.c"
+ DEPS
+ iree::base
+ iree::base::internal::flatcc::parsing
+ iree::schemas::executable_debug_info_c_fbs
+ PUBLIC
+)
+
+iree_cc_library(
+ NAME
file_cache
HDRS
"file_cache.h"
diff --git a/runtime/src/iree/hal/utils/executable_debug_info.c b/runtime/src/iree/hal/utils/executable_debug_info.c
new file mode 100644
index 0000000..0cd2149
--- /dev/null
+++ b/runtime/src/iree/hal/utils/executable_debug_info.c
@@ -0,0 +1,74 @@
+// Copyright 2024 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#include "iree/hal/utils/executable_debug_info.h"
+
+static iree_status_t iree_hal_debug_verify_string_nonempty(
+ const char* field_name, flatbuffers_string_t value) {
+ if (flatbuffers_string_len(value) == 0) {
+ return iree_make_status(
+ IREE_STATUS_INVALID_ARGUMENT,
+ "expected debug info field `%s` to contain a non-empty string value",
+ field_name);
+ }
+ return iree_ok_status();
+}
+
+static iree_status_t iree_hal_debug_verify_FileLineLocDef(
+ iree_hal_debug_FileLineLocDef_table_t def) {
+ if (!def) return iree_ok_status();
+ return iree_hal_debug_verify_string_nonempty(
+ "filename", iree_hal_debug_FileLineLocDef_filename_get(def));
+}
+
+iree_status_t iree_hal_debug_verify_export_def(
+ iree_hal_debug_ExportDef_table_t export_def) {
+ if (!export_def) return iree_ok_status();
+
+ IREE_RETURN_IF_ERROR(iree_hal_debug_verify_FileLineLocDef(
+ iree_hal_debug_ExportDef_location_get(export_def)));
+
+ iree_hal_debug_StageLocationDef_vec_t stage_locations_vec =
+ iree_hal_debug_ExportDef_stage_locations_get(export_def);
+ for (iree_host_size_t i = 0;
+ i < iree_hal_debug_StageLocationDef_vec_len(stage_locations_vec); ++i) {
+ iree_hal_debug_StageLocationDef_table_t stage_location_def =
+ iree_hal_debug_StageLocationDef_vec_at(stage_locations_vec, i);
+ if (!stage_location_def) {
+ return iree_make_status(IREE_STATUS_INVALID_ARGUMENT,
+ "stage_locations[%" PRIhsz "] has NULL value", i);
+ }
+ IREE_RETURN_IF_ERROR(iree_hal_debug_verify_string_nonempty(
+ "stage", iree_hal_debug_StageLocationDef_stage_get(
+ stage_location_def)),
+ "verifying stage_locations[%" PRIhsz "]", i);
+ IREE_RETURN_IF_ERROR(
+ iree_hal_debug_verify_FileLineLocDef(
+ iree_hal_debug_StageLocationDef_location_get(stage_location_def)),
+ "verifying stage_locations[%" PRIhsz "]", i);
+ }
+
+ return iree_ok_status();
+}
+
+void iree_hal_debug_publish_source_files(
+ iree_hal_debug_SourceFileDef_vec_t source_files_vec) {
+ if (!source_files_vec) return;
+#if IREE_TRACING_FEATURES & IREE_TRACING_FEATURE_INSTRUMENTATION
+ for (iree_host_size_t i = 0;
+ i < iree_hal_debug_SourceFileDef_vec_len(source_files_vec); ++i) {
+ iree_hal_debug_SourceFileDef_table_t source_file =
+ iree_hal_debug_SourceFileDef_vec_at(source_files_vec, i);
+ if (!source_file) continue;
+ flatbuffers_string_t path =
+ iree_hal_debug_SourceFileDef_path_get(source_file);
+ flatbuffers_uint8_vec_t content =
+ iree_hal_debug_SourceFileDef_content_get(source_file);
+ IREE_TRACE_PUBLISH_SOURCE_FILE(path, flatbuffers_string_len(path), content,
+ flatbuffers_uint8_vec_len(content));
+ }
+#endif // IREE_TRACING_FEATURES & IREE_TRACING_FEATURE_INSTRUMENTATION
+}
diff --git a/runtime/src/iree/hal/utils/executable_debug_info.h b/runtime/src/iree/hal/utils/executable_debug_info.h
new file mode 100644
index 0000000..bae9961
--- /dev/null
+++ b/runtime/src/iree/hal/utils/executable_debug_info.h
@@ -0,0 +1,36 @@
+// Copyright 2024 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#ifndef IREE_HAL_UTILS_EXECUTABLE_DEBUG_INFO_H_
+#define IREE_HAL_UTILS_EXECUTABLE_DEBUG_INFO_H_
+
+#include "iree/base/api.h"
+
+// flatcc schemas:
+#include "iree/base/internal/flatcc/parsing.h"
+#include "iree/schemas/executable_debug_info_reader.h"
+#include "iree/schemas/executable_debug_info_verifier.h"
+
+#ifdef __cplusplus
+extern "C" {
+#endif // __cplusplus
+
+// Verifies per-export debug info is valid.
+// Executables using debug info must call this as part of their verification.
+iree_status_t iree_hal_debug_verify_export_def(
+ iree_hal_debug_ExportDef_table_t export_def);
+
+// Publishes the given source files to any attached debug/trace providers.
+// This must be called prior to emitting any debug/trace events that reference
+// the files that are contained within.
+void iree_hal_debug_publish_source_files(
+ iree_hal_debug_SourceFileDef_vec_t source_files_vec);
+
+#ifdef __cplusplus
+} // extern "C"
+#endif // __cplusplus
+
+#endif // IREE_HAL_UTILS_EXECUTABLE_DEBUG_INFO_H_
diff --git a/runtime/src/iree/schemas/BUILD.bazel b/runtime/src/iree/schemas/BUILD.bazel
index 294c793..2f0959b 100644
--- a/runtime/src/iree/schemas/BUILD.bazel
+++ b/runtime/src/iree/schemas/BUILD.bazel
@@ -30,11 +30,12 @@
name = "cuda_executable_def_c_fbs",
srcs = ["cuda_executable_def.fbs"],
flatcc_args = FLATCC_ARGS,
+ includes = ["executable_debug_info.fbs"],
)
iree_flatbuffer_c_library(
- name = "rocm_executable_def_c_fbs",
- srcs = ["rocm_executable_def.fbs"],
+ name = "executable_debug_info_c_fbs",
+ srcs = ["executable_debug_info.fbs"],
flatcc_args = FLATCC_ARGS,
)
@@ -42,25 +43,38 @@
name = "metal_executable_def_c_fbs",
srcs = ["metal_executable_def.fbs"],
flatcc_args = FLATCC_ARGS,
+ includes = ["executable_debug_info.fbs"],
+)
+
+iree_flatbuffer_c_library(
+ name = "rocm_executable_def_c_fbs",
+ srcs = ["rocm_executable_def.fbs"],
+ flatcc_args = FLATCC_ARGS,
+ includes = ["executable_debug_info.fbs"],
)
iree_flatbuffer_c_library(
name = "spirv_executable_def_c_fbs",
srcs = ["spirv_executable_def.fbs"],
flatcc_args = FLATCC_ARGS,
+ includes = ["executable_debug_info.fbs"],
)
iree_flatbuffer_c_library(
name = "wgsl_executable_def_c_fbs",
srcs = ["wgsl_executable_def.fbs"],
flatcc_args = FLATCC_ARGS,
+ includes = ["executable_debug_info.fbs"],
)
iree_build_test(
name = "schema_build_test",
targets = [
":bytecode_module_def_c_fbs",
+ ":cuda_executable_def_c_fbs",
+ ":executable_debug_info_c_fbs",
":metal_executable_def_c_fbs",
+ ":rocm_executable_def_c_fbs",
":spirv_executable_def_c_fbs",
":wgsl_executable_def_c_fbs",
],
diff --git a/runtime/src/iree/schemas/CMakeLists.txt b/runtime/src/iree/schemas/CMakeLists.txt
index 776616e..cfbb850 100644
--- a/runtime/src/iree/schemas/CMakeLists.txt
+++ b/runtime/src/iree/schemas/CMakeLists.txt
@@ -33,14 +33,16 @@
"--builder"
"--verifier"
"--json"
+ INCLUDES
+ "executable_debug_info.fbs"
PUBLIC
)
flatbuffer_c_library(
NAME
- rocm_executable_def_c_fbs
+ executable_debug_info_c_fbs
SRCS
- "rocm_executable_def.fbs"
+ "executable_debug_info.fbs"
FLATCC_ARGS
"--reader"
"--builder"
@@ -59,6 +61,23 @@
"--builder"
"--verifier"
"--json"
+ INCLUDES
+ "executable_debug_info.fbs"
+ PUBLIC
+)
+
+flatbuffer_c_library(
+ NAME
+ rocm_executable_def_c_fbs
+ SRCS
+ "rocm_executable_def.fbs"
+ FLATCC_ARGS
+ "--reader"
+ "--builder"
+ "--verifier"
+ "--json"
+ INCLUDES
+ "executable_debug_info.fbs"
PUBLIC
)
@@ -72,6 +91,8 @@
"--builder"
"--verifier"
"--json"
+ INCLUDES
+ "executable_debug_info.fbs"
PUBLIC
)
@@ -85,6 +106,8 @@
"--builder"
"--verifier"
"--json"
+ INCLUDES
+ "executable_debug_info.fbs"
PUBLIC
)
diff --git a/runtime/src/iree/schemas/cuda_executable_def.fbs b/runtime/src/iree/schemas/cuda_executable_def.fbs
index b6713d4..0abc40a 100644
--- a/runtime/src/iree/schemas/cuda_executable_def.fbs
+++ b/runtime/src/iree/schemas/cuda_executable_def.fbs
@@ -4,6 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+include "iree/schemas/executable_debug_info.fbs";
+
namespace iree.hal.cuda;
// 'CUDA Executable'.
@@ -17,12 +19,6 @@
z:uint32;
}
-// Source code location denoted by a file name and line within that file.
-table FileLineLocDef {
- filename:string;
- line:int32;
-}
-
table ExecutableDef {
// A map of entry point ordinals to string names as used in the shader
// library.
@@ -39,12 +35,13 @@
// PTX string of the module.
ptx_image:string;
- // TODO(thomasraoux): Add potential cuBin binary specialized for some targets.
-
// A map of entry point ordinals to source locations.
// This information is optional and may be used by debuggers and profilers to
// associate executable entry points with the source that generated them.
- source_locations:[FileLineLocDef];
+ source_locations:[iree.hal.debug.FileLineLocDef];
+
+ // Embedded source files sorted ascending by path.
+ source_files:[iree.hal.debug.SourceFileDef];
}
root_type ExecutableDef;
diff --git a/runtime/src/iree/schemas/executable_debug_info.fbs b/runtime/src/iree/schemas/executable_debug_info.fbs
new file mode 100644
index 0000000..c3bceaa
--- /dev/null
+++ b/runtime/src/iree/schemas/executable_debug_info.fbs
@@ -0,0 +1,44 @@
+// Copyright 2024 The IREE Authors
+//
+// Licensed under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+namespace iree.hal.debug;
+
+// Source code location denoted by a file name and line within that file.
+table FileLineLocDef {
+ filename:string;
+ line:int32;
+}
+
+// Source location keyed by a string compilation stage name.
+table StageLocationDef {
+ stage:string;
+ location:FileLineLocDef;
+}
+
+// TODO(#18154): remove this when using ExportDef.
+// Table of stage locations sorted in ascending order by stage name.
+table StageLocationsDef {
+ locations:[StageLocationDef];
+}
+
+// Debug information for an exported function.
+// Empty/omitted if the compilation debug level is 0.
+table ExportDef {
+ // Source location in the canonical form to be presented in most tooling.
+ // Generally included with compilation debug level >= 1.
+ location:FileLineLocDef;
+
+ // Table of source locations keyed by compilation stage name.
+ // Sorted ascending by stage name.
+ // Generally included with compilation debug level >= 3.
+ stage_locations:[StageLocationDef];
+}
+
+// An embedded source file referenced by locations in the file.
+table SourceFileDef {
+ path:string;
+ content:[uint8];
+}
diff --git a/runtime/src/iree/schemas/metal_executable_def.fbs b/runtime/src/iree/schemas/metal_executable_def.fbs
index dc72781..fd03307 100644
--- a/runtime/src/iree/schemas/metal_executable_def.fbs
+++ b/runtime/src/iree/schemas/metal_executable_def.fbs
@@ -4,6 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+include "iree/schemas/executable_debug_info.fbs";
+
namespace iree.hal.metal;
// 'Metal Executable'.
@@ -42,6 +44,9 @@
shader_libraries:[string];
// Original Metal shader source code.
shader_sources:[string];
+
+ // Embedded source files sorted ascending by path.
+ source_files:[iree.hal.debug.SourceFileDef];
}
root_type ExecutableDef;
diff --git a/runtime/src/iree/schemas/rocm_executable_def.fbs b/runtime/src/iree/schemas/rocm_executable_def.fbs
index f368e1f..6781115 100644
--- a/runtime/src/iree/schemas/rocm_executable_def.fbs
+++ b/runtime/src/iree/schemas/rocm_executable_def.fbs
@@ -4,6 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+include "iree/schemas/executable_debug_info.fbs";
+
namespace iree.hal.rocm;
// 'ROCM Executable'.
@@ -17,30 +19,6 @@
z:uint32;
}
-// A struct for a source code location that consists of a file name and
-// a line number within that file.
-table FileLineLocDef {
- filename:string;
- line:int32;
-}
-
-// Source location keyed by a string compilation stage name.
-table StageLocationDef {
- stage:string;
- location:FileLineLocDef;
-}
-
-// Table of stage locations sorted in ascending order by stage name.
-table StageLocationsDef {
- locations:[StageLocationDef];
-}
-
-// An embedded source file referenced by locations in the file.
-table SourceFileDef {
- path:string;
- content:[uint8];
-}
-
table ExecutableDef {
// A map of entry point ordinals to string names as used in the shader
// library.
@@ -60,14 +38,14 @@
// A map of entry point ordinals to source locations.
// This information is optional and may be used by debuggers and profilers to
// associate executable entry points with the source that generated them.
- source_locations:[FileLineLocDef];
+ source_locations:[iree.hal.debug.FileLineLocDef];
// Table of source locations per entry point keyed by a string compilation
// stage name. Sorted ascending by name.
- stage_locations:[StageLocationsDef];
+ stage_locations:[iree.hal.debug.StageLocationsDef];
// Embedded source files sorted ascending by path.
- source_files:[SourceFileDef];
+ source_files:[iree.hal.debug.SourceFileDef];
}
root_type ExecutableDef;
diff --git a/runtime/src/iree/schemas/spirv_executable_def.fbs b/runtime/src/iree/schemas/spirv_executable_def.fbs
index 4eaea8f..a5aa17e 100644
--- a/runtime/src/iree/schemas/spirv_executable_def.fbs
+++ b/runtime/src/iree/schemas/spirv_executable_def.fbs
@@ -4,6 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+include "iree/schemas/executable_debug_info.fbs";
+
namespace iree.hal.spirv;
// 'SPIR-V Executable'.
@@ -15,29 +17,6 @@
code:[uint32];
}
-// Source code location denoted by a file name and line within that file.
-table FileLineLocDef {
- filename:string;
- line:int32;
-}
-
-// Source location keyed by a string compilation stage name.
-table StageLocationDef {
- stage:string;
- location:FileLineLocDef;
-}
-
-// Table of stage locations sorted in ascending order by stage name.
-table StageLocationsDef {
- locations:[StageLocationDef];
-}
-
-// An embedded source file referenced by locations in the file.
-table SourceFileDef {
- path:string;
- content:[uint8];
-}
-
// A SPIR-V shader module and runtime pipeline layout description.
// This information is used to create the VkShaderModule, VkPipelineLayout, and
// any required VkDescriptorSetLayouts.
@@ -63,14 +42,14 @@
// A map of entry point ordinals to source locations.
// This information is optional and may be used by debuggers and profilers to
// associate executable entry points with the source that generated them.
- source_locations:[FileLineLocDef];
+ source_locations:[iree.hal.debug.FileLineLocDef];
// Table of source locations per entry point keyed by a string compilation
// stage name. Sorted ascending by name.
- stage_locations:[StageLocationsDef];
+ stage_locations:[iree.hal.debug.StageLocationsDef];
// Embedded source files sorted ascending by path.
- source_files:[SourceFileDef];
+ source_files:[iree.hal.debug.SourceFileDef];
}
root_type ExecutableDef;
diff --git a/runtime/src/iree/schemas/wgsl_executable_def.fbs b/runtime/src/iree/schemas/wgsl_executable_def.fbs
index 79c821f..bba8f4c 100644
--- a/runtime/src/iree/schemas/wgsl_executable_def.fbs
+++ b/runtime/src/iree/schemas/wgsl_executable_def.fbs
@@ -4,6 +4,8 @@
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+include "iree/schemas/executable_debug_info.fbs";
+
namespace iree.hal.wgsl;
// 'WGSL Executable'.
@@ -28,6 +30,9 @@
// A mapping of executable entry point ordinals to the shader module in which
// they reside.
entry_points:[uint];
+
+ // Embedded source files sorted ascending by path.
+ source_files:[iree.hal.debug.SourceFileDef];
}
root_type ExecutableDef;
