From 3d5d65272d9be4678f242f918eb97e11362ff916 Mon Sep 17 00:00:00 2001 From: Ben Vanik Date: Tue, 13 Aug 2024 12:17:18 -0700 Subject: [PATCH 1/6] Fixing some incorrect TODOs referencing #18154. --- compiler/plugins/target/LLVMCPU/LLVMCPUTarget.cpp | 2 +- compiler/plugins/target/VMVX/VMVXTarget.cpp | 2 +- runtime/src/iree/hal/command_buffer.h | 8 ++++---- runtime/src/iree/hal/drivers/cuda/graph_command_buffer.c | 2 +- runtime/src/iree/hal/drivers/cuda/native_executable.c | 2 +- runtime/src/iree/hal/drivers/cuda/native_executable.h | 4 ++-- runtime/src/iree/hal/drivers/cuda/stream_command_buffer.c | 2 +- runtime/src/iree/hal/drivers/hip/graph_command_buffer.c | 2 +- runtime/src/iree/hal/drivers/hip/native_executable.c | 2 +- runtime/src/iree/hal/drivers/hip/native_executable.h | 4 ++-- runtime/src/iree/hal/drivers/hip/stream_command_buffer.c | 2 +- .../src/iree/hal/drivers/local_task/task_command_buffer.c | 4 ++-- runtime/src/iree/hal/local/executable_library.h | 2 +- runtime/src/iree/hal/local/inline_command_buffer.c | 4 ++-- 14 files changed, 21 insertions(+), 21 deletions(-) diff --git a/compiler/plugins/target/LLVMCPU/LLVMCPUTarget.cpp b/compiler/plugins/target/LLVMCPU/LLVMCPUTarget.cpp index 3a6337e766a5..78ffec60247c 100644 --- a/compiler/plugins/target/LLVMCPU/LLVMCPUTarget.cpp +++ b/compiler/plugins/target/LLVMCPU/LLVMCPUTarget.cpp @@ -398,7 +398,7 @@ class LLVMCPUTargetBackend final : public TargetBackend { // Specify the constant and binding information used to validate // dispatches. - // TODO(#18189): pack per-binding information bitfields. + // TODO(#18154): pack per-binding information bitfields. dispatchAttrs.constantCount = exportOp.getLayout().getPushConstants(); dispatchAttrs.bindingCount = exportOp.getLayout().getSetLayout(0).getBindings().size(); diff --git a/compiler/plugins/target/VMVX/VMVXTarget.cpp b/compiler/plugins/target/VMVX/VMVXTarget.cpp index 831eb8cd66a4..bef9f162d15b 100644 --- a/compiler/plugins/target/VMVX/VMVXTarget.cpp +++ b/compiler/plugins/target/VMVX/VMVXTarget.cpp @@ -132,7 +132,7 @@ class VMVXTargetBackend final : public TargetBackend { // Specify the constant and binding information used to validate // dispatches. - // TODO(#18189): pack per-binding information bitfields. + // TODO(#18154): pack per-binding information bitfields. if (auto layoutAttr = exportOp.getLayout()) { int64_t constantCount = layoutAttr.getPushConstants(); if (constantCount > 0) { diff --git a/runtime/src/iree/hal/command_buffer.h b/runtime/src/iree/hal/command_buffer.h index 43a876feef52..3b7fb77561af 100644 --- a/runtime/src/iree/hal/command_buffer.h +++ b/runtime/src/iree/hal/command_buffer.h @@ -766,8 +766,8 @@ IREE_API_EXPORT iree_status_t iree_hal_command_buffer_push_descriptor_set( // owning this queue. It must not be unregistered until all requests that use // it have completed. // -// Fails if the queue does not support dispatch operations (as indicated by -// can_dispatch). +// Fails if the queue does not support dispatch operations or +// IREE_HAL_COMMAND_CATEGORY_DISPATCH was not set. IREE_API_EXPORT iree_status_t iree_hal_command_buffer_dispatch( iree_hal_command_buffer_t* command_buffer, iree_hal_executable_t* executable, int32_t entry_point, @@ -799,8 +799,8 @@ IREE_API_EXPORT iree_status_t iree_hal_command_buffer_dispatch_indirect( // The provided constant data and binding list will be recorded into the command // buffer and need not remain live beyond the call. // -// Fails if the queue does not support dispatch operations (as indicated by -// can_dispatch). +// Fails if the queue does not support dispatch operations or +// IREE_HAL_COMMAND_CATEGORY_DISPATCH was not set. IREE_API_EXPORT iree_status_t iree_hal_command_buffer_dispatch2( iree_hal_command_buffer_t* command_buffer, iree_hal_executable_t* executable, int32_t entry_point, diff --git a/runtime/src/iree/hal/drivers/cuda/graph_command_buffer.c b/runtime/src/iree/hal/drivers/cuda/graph_command_buffer.c index 68d4d34668bb..276879f5d498 100644 --- a/runtime/src/iree/hal/drivers/cuda/graph_command_buffer.c +++ b/runtime/src/iree/hal/drivers/cuda/graph_command_buffer.c @@ -59,7 +59,7 @@ typedef struct iree_hal_cuda_graph_command_buffer_t { // Iteratively constructed batch of collective operations. iree_hal_collective_batch_t collective_batch; - // TODO(#18189): drop state used by legacy bindings mechanism. + // TODO(#18154): drop state used by legacy bindings mechanism. int32_t push_constants[IREE_HAL_CUDA_MAX_PUSH_CONSTANT_COUNT]; struct { CUdeviceptr bindings[IREE_HAL_CUDA_MAX_DESCRIPTOR_SET_BINDING_COUNT]; diff --git a/runtime/src/iree/hal/drivers/cuda/native_executable.c b/runtime/src/iree/hal/drivers/cuda/native_executable.c index 06a7ffc1bbbe..d6958d8c9a96 100644 --- a/runtime/src/iree/hal/drivers/cuda/native_executable.c +++ b/runtime/src/iree/hal/drivers/cuda/native_executable.c @@ -224,7 +224,7 @@ iree_status_t iree_hal_cuda_native_executable_create( } if (!iree_status_is_ok(status)) break; - // TODO(#18189): embed all of this on a single flatbuffer table + // TODO(#18154): embed all of this on a single flatbuffer table // per-export. // // Package required parameters for kernel launches for each entry point. diff --git a/runtime/src/iree/hal/drivers/cuda/native_executable.h b/runtime/src/iree/hal/drivers/cuda/native_executable.h index 226cedaa6b44..3f6faf5ec8ef 100644 --- a/runtime/src/iree/hal/drivers/cuda/native_executable.h +++ b/runtime/src/iree/hal/drivers/cuda/native_executable.h @@ -20,12 +20,12 @@ extern "C" { #endif // __cplusplus typedef struct iree_hal_cuda_kernel_info_t { - // TODO(#18189): remove when using simplified bindings. + // TODO(#18154): remove when using simplified bindings. iree_hal_pipeline_layout_t* layout; CUfunction function; uint32_t constant_count; uint32_t binding_count; - // TODO(#18189): add bitfield indicating indirect bindings. + // TODO(#18154): add bitfield indicating indirect bindings. uint32_t block_size[3]; uint32_t shared_memory_size; diff --git a/runtime/src/iree/hal/drivers/cuda/stream_command_buffer.c b/runtime/src/iree/hal/drivers/cuda/stream_command_buffer.c index a9b50fc19f4a..908ae10b1635 100644 --- a/runtime/src/iree/hal/drivers/cuda/stream_command_buffer.c +++ b/runtime/src/iree/hal/drivers/cuda/stream_command_buffer.c @@ -39,7 +39,7 @@ typedef struct iree_hal_cuda_stream_command_buffer_t { // Iteratively constructed batch of collective operations. iree_hal_collective_batch_t collective_batch; - // TODO(#18189): drop state used by legacy bindings mechanism. + // TODO(#18154): drop state used by legacy bindings mechanism. int32_t push_constants[IREE_HAL_CUDA_MAX_PUSH_CONSTANT_COUNT]; struct { CUdeviceptr bindings[IREE_HAL_CUDA_MAX_DESCRIPTOR_SET_BINDING_COUNT]; diff --git a/runtime/src/iree/hal/drivers/hip/graph_command_buffer.c b/runtime/src/iree/hal/drivers/hip/graph_command_buffer.c index 99b3538caf77..88bdf841334c 100644 --- a/runtime/src/iree/hal/drivers/hip/graph_command_buffer.c +++ b/runtime/src/iree/hal/drivers/hip/graph_command_buffer.c @@ -60,7 +60,7 @@ typedef struct iree_hal_hip_graph_command_buffer_t { // Iteratively constructed batch of collective operations. iree_hal_collective_batch_t collective_batch; - // TODO(#18189): drop state used by legacy bindings mechanism. + // TODO(#18154): drop state used by legacy bindings mechanism. int32_t push_constants[IREE_HAL_HIP_MAX_PUSH_CONSTANT_COUNT]; struct { hipDeviceptr_t bindings[IREE_HAL_HIP_MAX_DESCRIPTOR_SET_BINDING_COUNT]; diff --git a/runtime/src/iree/hal/drivers/hip/native_executable.c b/runtime/src/iree/hal/drivers/hip/native_executable.c index 19caae90aa22..5e06e81116f2 100644 --- a/runtime/src/iree/hal/drivers/hip/native_executable.c +++ b/runtime/src/iree/hal/drivers/hip/native_executable.c @@ -243,7 +243,7 @@ iree_status_t iree_hal_hip_native_executable_create( } if (!iree_status_is_ok(status)) break; - // TODO(#18189): embed all of this on a single flatbuffer table + // TODO(#18154): embed all of this on a single flatbuffer table // per-export. // // Package required parameters for kernel launches for each entry point. diff --git a/runtime/src/iree/hal/drivers/hip/native_executable.h b/runtime/src/iree/hal/drivers/hip/native_executable.h index d2b1a319de5c..d68880a76cab 100644 --- a/runtime/src/iree/hal/drivers/hip/native_executable.h +++ b/runtime/src/iree/hal/drivers/hip/native_executable.h @@ -20,12 +20,12 @@ extern "C" { #endif // __cplusplus typedef struct iree_hal_hip_kernel_info_t { - // TODO(#18189): remove when using simplified bindings. + // TODO(#18154): remove when using simplified bindings. iree_hal_pipeline_layout_t* layout; hipFunction_t function; uint32_t constant_count; uint32_t binding_count; - // TODO(#18189): add bitfield indicating indirect bindings. + // TODO(#18154): add bitfield indicating indirect bindings. uint32_t block_size[3]; uint32_t shared_memory_size; diff --git a/runtime/src/iree/hal/drivers/hip/stream_command_buffer.c b/runtime/src/iree/hal/drivers/hip/stream_command_buffer.c index e4ffac2200a9..9b92e2317083 100644 --- a/runtime/src/iree/hal/drivers/hip/stream_command_buffer.c +++ b/runtime/src/iree/hal/drivers/hip/stream_command_buffer.c @@ -41,7 +41,7 @@ typedef struct iree_hal_hip_stream_command_buffer_t { // Iteratively constructed batch of collective operations. iree_hal_collective_batch_t collective_batch; - // TODO(#18189): drop state used by legacy bindings mechanism. + // TODO(#18154): drop state used by legacy bindings mechanism. int32_t push_constants[IREE_HAL_HIP_MAX_PUSH_CONSTANT_COUNT]; struct { hipDeviceptr_t bindings[IREE_HAL_HIP_MAX_DESCRIPTOR_SET_BINDING_COUNT]; diff --git a/runtime/src/iree/hal/drivers/local_task/task_command_buffer.c b/runtime/src/iree/hal/drivers/local_task/task_command_buffer.c index 3b0a9ba18617..1fbeb1779827 100644 --- a/runtime/src/iree/hal/drivers/local_task/task_command_buffer.c +++ b/runtime/src/iree/hal/drivers/local_task/task_command_buffer.c @@ -78,7 +78,7 @@ typedef struct iree_hal_task_command_buffer_t { // All execution tasks emitted that must execute after |open_barrier|. iree_task_list_t open_tasks; - // TODO(#18189): remove legacy binding state. + // TODO(#18154): remove legacy binding state. // A flattened list of all available descriptor set bindings. // As descriptor sets are pushed/bound the bindings will be updated to // represent the fully-translated binding data pointer. @@ -90,7 +90,7 @@ typedef struct iree_hal_task_command_buffer_t { binding_lengths[IREE_HAL_LOCAL_MAX_DESCRIPTOR_SET_COUNT * IREE_HAL_LOCAL_MAX_DESCRIPTOR_BINDING_COUNT]; - // TODO(#18189): remove legacy push constant state. + // TODO(#18154): remove legacy push constant state. // All available push constants updated each time push_constants is called. // Reset only with the command buffer and otherwise will maintain its values // during recording to allow for partial push_constants updates. diff --git a/runtime/src/iree/hal/local/executable_library.h b/runtime/src/iree/hal/local/executable_library.h index d45b477d5573..f7883ba1ce78 100644 --- a/runtime/src/iree/hal/local/executable_library.h +++ b/runtime/src/iree/hal/local/executable_library.h @@ -392,7 +392,7 @@ typedef struct iree_hal_executable_dispatch_attrs_v0_t { uint8_t constant_count; // Total number of bindings used by the dispatch. uint8_t binding_count; - // TODO(#18189): add ~8 uint64_t fields for binding bits (readonly/indirect). + // TODO(#18154): add ~8 uint64_t fields for binding bits (readonly/indirect). } iree_hal_executable_dispatch_attrs_v0_t; static_assert(sizeof(iree_hal_executable_dispatch_attrs_v0_t) == 4, "uint32_t"); diff --git a/runtime/src/iree/hal/local/inline_command_buffer.c b/runtime/src/iree/hal/local/inline_command_buffer.c index 2e0465c0bcf7..90d3641a8efc 100644 --- a/runtime/src/iree/hal/local/inline_command_buffer.c +++ b/runtime/src/iree/hal/local/inline_command_buffer.c @@ -28,7 +28,7 @@ typedef struct iree_hal_inline_command_buffer_t { iree_allocator_t host_allocator; struct { - // TODO(#18189): remove legacy bindings state. + // TODO(#18154): remove legacy bindings state. // // A flattened list of all available descriptor set bindings. // As descriptor sets are pushed/bound the bindings will be updated to @@ -38,7 +38,7 @@ typedef struct iree_hal_inline_command_buffer_t { size_t full_binding_lengths[IREE_HAL_LOCAL_MAX_DESCRIPTOR_SET_COUNT * IREE_HAL_LOCAL_MAX_DESCRIPTOR_BINDING_COUNT]; - // TODO(#18189): remove legacy push constant state. + // TODO(#18154): remove legacy push constant state. // // All available push constants updated each time push_constants is called. // Reset only with the command buffer and otherwise will maintain its values From 8794c92b25fe979df71355a011b3707bf2e2ea39 Mon Sep 17 00:00:00 2001 From: Ben Vanik Date: Tue, 13 Aug 2024 12:19:31 -0700 Subject: [PATCH 2/6] Fixing ROCM/CUDA BlockSizeDef -> BlockSize. (tables have Def, structs don't) --- compiler/plugins/target/CUDA/CUDATarget.cpp | 6 +++--- compiler/plugins/target/ROCM/ROCMTarget.cpp | 6 +++--- docs/website/docs/community/blog/posts/cuda-backend.md | 2 +- runtime/src/iree/hal/drivers/cuda/native_executable.c | 6 +++--- runtime/src/iree/hal/drivers/hip/native_executable.c | 6 +++--- runtime/src/iree/schemas/cuda_executable_def.fbs | 4 ++-- runtime/src/iree/schemas/rocm_executable_def.fbs | 4 ++-- 7 files changed, 17 insertions(+), 17 deletions(-) diff --git a/compiler/plugins/target/CUDA/CUDATarget.cpp b/compiler/plugins/target/CUDA/CUDATarget.cpp index 7e78ae3eeb2c..eb05e187fb03 100644 --- a/compiler/plugins/target/CUDA/CUDATarget.cpp +++ b/compiler/plugins/target/CUDA/CUDATarget.cpp @@ -665,12 +665,12 @@ class CUDATargetBackend final : public TargetBackend { std::string gpuImage = produceGpuImage(options, targetArch, ptxImage); auto gpuImageRef = flatbuffers_string_create(builder, gpuImage.c_str(), gpuImage.size()); - iree_hal_cuda_BlockSizeDef_vec_start(builder); + iree_hal_cuda_BlockSize_vec_start(builder); for (const auto &workgroupSize : workgroupSizes) { - iree_hal_cuda_BlockSizeDef_vec_push_create( + iree_hal_cuda_BlockSize_vec_push_create( builder, workgroupSize[0], workgroupSize[1], workgroupSize[2]); } - auto blockSizesRef = iree_hal_cuda_BlockSizeDef_vec_end(builder); + auto blockSizesRef = iree_hal_cuda_BlockSize_vec_end(builder); auto workgroupLocalMemoriesRef = builder.createInt32Vec(workgroupLocalMemories); auto entryPointsRef = builder.createStringVec(entryPointNames); diff --git a/compiler/plugins/target/ROCM/ROCMTarget.cpp b/compiler/plugins/target/ROCM/ROCMTarget.cpp index d6cc400a299e..ff9a3694f037 100644 --- a/compiler/plugins/target/ROCM/ROCMTarget.cpp +++ b/compiler/plugins/target/ROCM/ROCMTarget.cpp @@ -651,16 +651,16 @@ class ROCMTargetBackend final : public TargetBackend { targetHSACO.size()); auto entryPointsRef = builder.createStringVec(entryPointNames); - iree_hal_rocm_BlockSizeDef_vec_start(builder); + iree_hal_rocm_BlockSize_vec_start(builder); auto blockSizes = workgroupSizes.begin(); for (int i = 0, e = entryPointNames.size(); i < e; ++i) { - iree_hal_rocm_BlockSizeDef_vec_push_create( + iree_hal_rocm_BlockSize_vec_push_create( builder, (*blockSizes)[0], (*blockSizes)[1], (*blockSizes)[2]); ++blockSizes; } auto workgroupLocalMemoriesRef = builder.createInt32Vec(workgroupLocalMemories); - auto blockSizesRef = iree_hal_rocm_BlockSizeDef_vec_end(builder); + auto blockSizesRef = iree_hal_rocm_BlockSize_vec_end(builder); iree_hal_rocm_ExecutableDef_entry_points_add(builder, entryPointsRef); iree_hal_rocm_ExecutableDef_block_sizes_add(builder, blockSizesRef); iree_hal_rocm_ExecutableDef_shared_memory_sizes_add( diff --git a/docs/website/docs/community/blog/posts/cuda-backend.md b/docs/website/docs/community/blog/posts/cuda-backend.md index ab2ee219c4ff..7c5adc38adf7 100644 --- a/docs/website/docs/community/blog/posts/cuda-backend.md +++ b/docs/website/docs/community/blog/posts/cuda-backend.md @@ -82,7 +82,7 @@ table CUDAExecutableDef { entry_points:[string]; // Block sizes for each entry point. - block_sizes:[CUDABlockSizeDef]; + block_sizes:[CUDABlockSize]; // PTX string of the module. ptx_image:string; diff --git a/runtime/src/iree/hal/drivers/cuda/native_executable.c b/runtime/src/iree/hal/drivers/cuda/native_executable.c index d6958d8c9a96..18ab1e62e592 100644 --- a/runtime/src/iree/hal/drivers/cuda/native_executable.c +++ b/runtime/src/iree/hal/drivers/cuda/native_executable.c @@ -84,9 +84,9 @@ static iree_status_t iree_hal_cuda_native_executable_flatbuffer_verify( } } - iree_hal_cuda_BlockSizeDef_vec_t block_sizes_vec = + iree_hal_cuda_BlockSize_vec_t block_sizes_vec = iree_hal_cuda_ExecutableDef_block_sizes_get(executable_def); - size_t block_size_count = iree_hal_cuda_BlockSizeDef_vec_len(block_sizes_vec); + size_t block_size_count = iree_hal_cuda_BlockSize_vec_len(block_sizes_vec); if (block_size_count == 0) { return iree_make_status(IREE_STATUS_INVALID_ARGUMENT, "no block sizes present"); @@ -134,7 +134,7 @@ iree_status_t iree_hal_cuda_native_executable_create( iree_hal_cuda_ExecutableDef_shared_memory_size_get(executable_def); flatbuffers_string_vec_t entry_points_vec = iree_hal_cuda_ExecutableDef_entry_points_get(executable_def); - iree_hal_cuda_BlockSizeDef_vec_t block_sizes_vec = + iree_hal_cuda_BlockSize_vec_t block_sizes_vec = iree_hal_cuda_ExecutableDef_block_sizes_get(executable_def); iree_host_size_t entry_point_count = flatbuffers_string_vec_len(entry_points_vec); diff --git a/runtime/src/iree/hal/drivers/hip/native_executable.c b/runtime/src/iree/hal/drivers/hip/native_executable.c index 5e06e81116f2..5feab821e06d 100644 --- a/runtime/src/iree/hal/drivers/hip/native_executable.c +++ b/runtime/src/iree/hal/drivers/hip/native_executable.c @@ -88,9 +88,9 @@ static iree_status_t iree_hal_hip_native_executable_flatbuffer_verify( } } - iree_hal_rocm_BlockSizeDef_vec_t block_sizes_vec = + iree_hal_rocm_BlockSize_vec_t block_sizes_vec = iree_hal_rocm_ExecutableDef_block_sizes_get(executable_def); - size_t block_size_count = iree_hal_rocm_BlockSizeDef_vec_len(block_sizes_vec); + size_t block_size_count = iree_hal_rocm_BlockSize_vec_len(block_sizes_vec); if (entry_point_count != block_size_count) { return iree_make_status( IREE_STATUS_INVALID_ARGUMENT, @@ -141,7 +141,7 @@ iree_status_t iree_hal_hip_native_executable_create( flatbuffers_string_vec_t entry_points_vec = iree_hal_rocm_ExecutableDef_entry_points_get(executable_def); - iree_hal_rocm_BlockSizeDef_vec_t block_sizes_vec = + iree_hal_rocm_BlockSize_vec_t block_sizes_vec = iree_hal_rocm_ExecutableDef_block_sizes_get(executable_def); flatbuffers_uint32_vec_t shared_memory_sizes_vec = iree_hal_rocm_ExecutableDef_shared_memory_sizes_get(executable_def); diff --git a/runtime/src/iree/schemas/cuda_executable_def.fbs b/runtime/src/iree/schemas/cuda_executable_def.fbs index df78a7df8270..b6713d493d7a 100644 --- a/runtime/src/iree/schemas/cuda_executable_def.fbs +++ b/runtime/src/iree/schemas/cuda_executable_def.fbs @@ -11,7 +11,7 @@ file_identifier "CUDA"; file_extension "cuda"; // A struct for the kernel block size along each dimensions. -struct BlockSizeDef { +struct BlockSize { x:uint32; y:uint32; z:uint32; @@ -32,7 +32,7 @@ table ExecutableDef { // // Currently the thread group size/block size is decided during code gen but // in CUDA it is set by the runtime. - block_sizes:[BlockSizeDef]; + block_sizes:[BlockSize]; // Size of dynamic shared memory. shared_memory_size:[uint32]; diff --git a/runtime/src/iree/schemas/rocm_executable_def.fbs b/runtime/src/iree/schemas/rocm_executable_def.fbs index 6df6d022de33..f368e1f43fda 100644 --- a/runtime/src/iree/schemas/rocm_executable_def.fbs +++ b/runtime/src/iree/schemas/rocm_executable_def.fbs @@ -11,7 +11,7 @@ file_identifier "ROCM"; file_extension "rocm"; // A struct for the kernel block size along each dimensions. -struct BlockSizeDef { +struct BlockSize { x:uint32; y:uint32; z:uint32; @@ -48,7 +48,7 @@ table ExecutableDef { // Block sizes for each entry point. // This list has the same size as the entry_points list. - block_sizes:[BlockSizeDef]; + block_sizes:[BlockSize]; // Size of dynamic shared memory. // This list has the same size as the entry_points list. From 39fd06f4fd50dffb3d7144a94d302895a262b9f5 Mon Sep 17 00:00:00 2001 From: Ben Vanik Date: Tue, 13 Aug 2024 12:22:10 -0700 Subject: [PATCH 3/6] Renaming Metal kernel_library -> executable. It's a HAL executable and will contain multiple kernel libraries. --- .../design-docs/metal-hal-driver.md | 4 +- .../src/iree/hal/drivers/metal/CMakeLists.txt | 4 +- .../hal/drivers/metal/builtin_executables.h | 2 +- .../hal/drivers/metal/builtin_executables.m | 9 ++-- .../hal/drivers/metal/direct_command_buffer.m | 6 +-- .../metal/{kernel_library.h => executable.h} | 12 ++--- .../metal/{kernel_library.m => executable.m} | 54 +++++++++---------- .../hal/drivers/metal/nop_executable_cache.m | 6 +-- 8 files changed, 47 insertions(+), 50 deletions(-) rename runtime/src/iree/hal/drivers/metal/{kernel_library.h => executable.h} (87%) rename runtime/src/iree/hal/drivers/metal/{kernel_library.m => executable.m} (89%) diff --git a/docs/website/docs/developers/design-docs/metal-hal-driver.md b/docs/website/docs/developers/design-docs/metal-hal-driver.md index 855ddad095f5..d0b6c75ce454 100644 --- a/docs/website/docs/developers/design-docs/metal-hal-driver.md +++ b/docs/website/docs/developers/design-docs/metal-hal-driver.md @@ -195,7 +195,7 @@ IREE [`iree_hal_buffer_t`][hal-buffer] maps Metal `MTLBuffer`. See IREE [`iree_hal_executable_t`][hal-executable] represents a GPU program archive with a driver-defined format. It maps naturally to Metal [`MTLLibrary`][mtl-library]. An entry point in a `MTLLibrary` is a [`MTLFunction`][mtl-function]. We define -[`iree_hal_metal_kernel_params_t`][metal-kernel-library] to wrap around a +[`iree_hal_metal_executable_t`][metal-executable] to wrap around a `MTLLibrary`, its `MTLFunction`s, and also `MTLComputePipelineState` objects constructed from `MTLFunction`s. @@ -328,7 +328,7 @@ with the current active `MTLComputeCommandEncoder`: [hal-semaphore]: https://github.com/iree-org/iree/blob/main/runtime/src/iree/hal/semaphore.h [metal-device]: https://github.com/iree-org/iree/tree/main/experimental/metal/metal_device.h [metal-driver]: https://github.com/iree-org/iree/tree/main/experimental/metal/metal_driver.h -[metal-kernel-library]: https://github.com/iree-org/iree/tree/main/experimental/metal/kernel_library.h +[metal-executable]: https://github.com/iree-org/iree/tree/main/experimental/metal/executable.h [metal-shared-event]: https://github.com/iree-org/iree/tree/main/experimental/metal/shared_event.h [metal-spirv-target]: https://github.com/iree-org/iree/tree/main/compiler/plugins/target/MetalSPIRV [metal-builtin-kernels]: https://github.com/iree-org/iree/tree/main/runtime/src/iree/hal/drivers/metal/builtin/ diff --git a/runtime/src/iree/hal/drivers/metal/CMakeLists.txt b/runtime/src/iree/hal/drivers/metal/CMakeLists.txt index c3186a109824..85356f30ec91 100644 --- a/runtime/src/iree/hal/drivers/metal/CMakeLists.txt +++ b/runtime/src/iree/hal/drivers/metal/CMakeLists.txt @@ -19,8 +19,8 @@ iree_cc_library( "direct_allocator.m" "direct_command_buffer.h" "direct_command_buffer.m" - "kernel_library.h" - "kernel_library.m" + "executable.h" + "executable.m" "metal_buffer.h" "metal_buffer.m" "metal_device.m" diff --git a/runtime/src/iree/hal/drivers/metal/builtin_executables.h b/runtime/src/iree/hal/drivers/metal/builtin_executables.h index 08fc065e3811..ab74cb3d2f63 100644 --- a/runtime/src/iree/hal/drivers/metal/builtin_executables.h +++ b/runtime/src/iree/hal/drivers/metal/builtin_executables.h @@ -11,7 +11,7 @@ #include "iree/base/api.h" #include "iree/hal/api.h" -#include "iree/hal/drivers/metal/kernel_library.h" +#include "iree/hal/drivers/metal/executable.h" #ifdef __cplusplus extern "C" { diff --git a/runtime/src/iree/hal/drivers/metal/builtin_executables.m b/runtime/src/iree/hal/drivers/metal/builtin_executables.m index 77912ab9afee..faf5045a90bd 100644 --- a/runtime/src/iree/hal/drivers/metal/builtin_executables.m +++ b/runtime/src/iree/hal/drivers/metal/builtin_executables.m @@ -83,16 +83,13 @@ iree_status_t iree_hal_metal_builtin_executable_create( if (!iree_status_is_ok(status)) break; // Package required parameters for kernel launches for each entry point. + // Thread group size for builtin executables are determined at runtime dispatch time. + // We don't need the layout information for builtins either. iree_hal_metal_kernel_params_t* params = &executable->entry_points[i]; + memset(params, 0, sizeof(*params)); params->library = library; params->function = function; params->pso = pso; - // Thread group size for builtin executables are determined at runtime dispatch time. - params->threadgroup_size[0] = 0; - params->threadgroup_size[1] = 0; - params->threadgroup_size[2] = 0; - // We don't need the layout parameter for builtin executables too. - params->layout = NULL; // Stash the entry point name in the string table for use when tracing. IREE_TRACE({ params->function_name = IREE_SV(entry_point); }); diff --git a/runtime/src/iree/hal/drivers/metal/direct_command_buffer.m b/runtime/src/iree/hal/drivers/metal/direct_command_buffer.m index fbf6374e7fcc..6754c12b3ab7 100644 --- a/runtime/src/iree/hal/drivers/metal/direct_command_buffer.m +++ b/runtime/src/iree/hal/drivers/metal/direct_command_buffer.m @@ -13,7 +13,7 @@ #include "iree/base/tracing.h" #include "iree/hal/api.h" #include "iree/hal/drivers/metal/builtin_executables.h" -#include "iree/hal/drivers/metal/kernel_library.h" +#include "iree/hal/drivers/metal/executable.h" #include "iree/hal/drivers/metal/metal_buffer.h" #include "iree/hal/drivers/metal/metal_device.h" #include "iree/hal/drivers/metal/pipeline_layout.h" @@ -962,7 +962,7 @@ static iree_status_t iree_hal_metal_command_segment_create_dispatch( z0, iree_hal_resource_set_insert(command_buffer->resource_set, 1, &executable)); iree_hal_metal_kernel_params_t kernel_params; - IREE_RETURN_AND_END_ZONE_IF_ERROR(z0, iree_hal_metal_kernel_library_entry_point_kernel_params( + IREE_RETURN_AND_END_ZONE_IF_ERROR(z0, iree_hal_metal_executable_entry_point_kernel_params( executable, entry_point, &kernel_params)); // Allocate the command segment and keep track of all necessary API data. @@ -1144,7 +1144,7 @@ static iree_status_t iree_hal_metal_command_segment_create_dispatch2( z0, iree_hal_resource_set_insert(command_buffer->resource_set, 1, &executable)); iree_hal_metal_kernel_params_t kernel_params; - IREE_RETURN_AND_END_ZONE_IF_ERROR(z0, iree_hal_metal_kernel_library_entry_point_kernel_params( + IREE_RETURN_AND_END_ZONE_IF_ERROR(z0, iree_hal_metal_executable_entry_point_kernel_params( executable, entry_point, &kernel_params)); // Allocate the command segment and keep track of all necessary API data. diff --git a/runtime/src/iree/hal/drivers/metal/kernel_library.h b/runtime/src/iree/hal/drivers/metal/executable.h similarity index 87% rename from runtime/src/iree/hal/drivers/metal/kernel_library.h rename to runtime/src/iree/hal/drivers/metal/executable.h index aa7c95742b77..701841f7b452 100644 --- a/runtime/src/iree/hal/drivers/metal/kernel_library.h +++ b/runtime/src/iree/hal/drivers/metal/executable.h @@ -4,14 +4,13 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#ifndef IREE_HAL_DRIVERS_METAL_KERNEL_LIBRARY_H_ -#define IREE_HAL_DRIVERS_METAL_KERNEL_LIBRARY_H_ +#ifndef IREE_HAL_DRIVERS_METAL_EXECUTABLE_H_ +#define IREE_HAL_DRIVERS_METAL_EXECUTABLE_H_ #import #include #include "iree/base/api.h" -#include "iree/base/tracing.h" #include "iree/hal/api.h" #ifdef __cplusplus @@ -24,6 +23,7 @@ typedef struct iree_hal_metal_kernel_params_t { id function; id pso; uint32_t threadgroup_size[3]; + // TODO(#18154): remove layout in simplified bindings. iree_hal_pipeline_layout_t* layout; IREE_TRACE(iree_string_view_t function_name;) } iree_hal_metal_kernel_params_t; @@ -40,12 +40,12 @@ typedef struct iree_hal_metal_kernel_params_t { // // |out_executable| must be released by the caller (see // iree_hal_executable_release). -iree_status_t iree_hal_metal_kernel_library_create( +iree_status_t iree_hal_metal_executable_create( id device, const iree_hal_executable_params_t* executable_params, iree_allocator_t host_allocator, iree_hal_executable_t** out_executable); // Returns the kernel launch parameters for the given |entry_point|. -iree_status_t iree_hal_metal_kernel_library_entry_point_kernel_params( +iree_status_t iree_hal_metal_executable_entry_point_kernel_params( const iree_hal_executable_t* executable, int32_t entry_point, iree_hal_metal_kernel_params_t* out_params); @@ -61,4 +61,4 @@ iree_status_t iree_hal_metal_compile_msl_and_create_pipeline_object( } // extern "C" #endif // __cplusplus -#endif // IREE_HAL_DRIVERS_METAL_KERNEL_LIBRARY_H_ +#endif // IREE_HAL_DRIVERS_METAL_EXECUTABLE_H_ diff --git a/runtime/src/iree/hal/drivers/metal/kernel_library.m b/runtime/src/iree/hal/drivers/metal/executable.m similarity index 89% rename from runtime/src/iree/hal/drivers/metal/kernel_library.m rename to runtime/src/iree/hal/drivers/metal/executable.m index f759f5e69c3b..b30ea157f056 100644 --- a/runtime/src/iree/hal/drivers/metal/kernel_library.m +++ b/runtime/src/iree/hal/drivers/metal/executable.m @@ -4,7 +4,7 @@ // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#include "iree/hal/drivers/metal/kernel_library.h" +#include "iree/hal/drivers/metal/executable.h" #include @@ -15,7 +15,7 @@ #include "iree/schemas/metal_executable_def_reader.h" #include "iree/schemas/metal_executable_def_verifier.h" -typedef struct iree_hal_metal_kernel_library_t { +typedef struct iree_hal_metal_executable_t { // Abstract resource used for injecting reference counting and vtable; must be at offset 0. iree_hal_resource_t resource; @@ -23,20 +23,20 @@ iree_host_size_t entry_point_count; iree_hal_metal_kernel_params_t entry_points[]; -} iree_hal_metal_kernel_library_t; +} iree_hal_metal_executable_t; -static const iree_hal_executable_vtable_t iree_hal_metal_kernel_library_vtable; +static const iree_hal_executable_vtable_t iree_hal_metal_executable_vtable; -static iree_hal_metal_kernel_library_t* iree_hal_metal_kernel_library_cast( +static iree_hal_metal_executable_t* iree_hal_metal_executable_cast( iree_hal_executable_t* base_value) { - IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_metal_kernel_library_vtable); - return (iree_hal_metal_kernel_library_t*)base_value; + IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_metal_executable_vtable); + return (iree_hal_metal_executable_t*)base_value; } -static const iree_hal_metal_kernel_library_t* iree_hal_metal_kernel_library_const_cast( +static const iree_hal_metal_executable_t* iree_hal_metal_executable_const_cast( const iree_hal_executable_t* base_value) { - IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_metal_kernel_library_vtable); - return (const iree_hal_metal_kernel_library_t*)base_value; + IREE_HAL_ASSERT_TYPE(base_value, &iree_hal_metal_executable_vtable); + return (const iree_hal_metal_executable_t*)base_value; } // Verifies the structure of the flatbuffer so that we can avoid doing so during runtime. @@ -44,7 +44,7 @@ // There are still some conditions we must be aware of (such as omitted names on functions with // internal linkage), however we shouldn't need to bounds check anything within the flatbuffer // after this succeeds. -static iree_status_t iree_hal_metal_kernel_library_flatbuffer_verify( +static iree_status_t iree_hal_metal_executable_flatbuffer_verify( iree_const_byte_span_t flatbuffer_data) { if (!flatbuffer_data.data || flatbuffer_data.data_length < 16) { return iree_make_status(IREE_STATUS_INVALID_ARGUMENT, @@ -192,7 +192,7 @@ static iree_status_t iree_hal_metal_load_mtllib(iree_const_byte_span_t source_da // Creates MTL compute pipeline objects for the given |entry_point| in |library| and writes to // |out_function| and |out_pso|. The caller should release |out_function| and |out_pso| after done. -static iree_status_t iree_hal_metal_create_pipline_object( +static iree_status_t iree_hal_metal_create_pipeline_object( id library, iree_string_view_t entry_point, const char* source_code, id device, id* out_function, id* out_pso) { @autoreleasepool { @@ -226,11 +226,11 @@ iree_status_t iree_hal_metal_compile_msl_and_create_pipeline_object( id* out_pso) { IREE_RETURN_IF_ERROR( iree_hal_metal_compile_msl(source_code, entry_point, device, compile_options, out_library)); - return iree_hal_metal_create_pipline_object(*out_library, entry_point, source_code.data, device, - out_function, out_pso); + return iree_hal_metal_create_pipeline_object(*out_library, entry_point, source_code.data, device, + out_function, out_pso); } -iree_status_t iree_hal_metal_kernel_library_create( +iree_status_t iree_hal_metal_executable_create( id device, const iree_hal_executable_params_t* executable_params, iree_allocator_t host_allocator, iree_hal_executable_t** out_executable) { IREE_ASSERT_ARGUMENT(executable_params); @@ -238,10 +238,10 @@ iree_status_t iree_hal_metal_kernel_library_create( *out_executable = NULL; IREE_TRACE_ZONE_BEGIN(z0); - iree_hal_metal_kernel_library_t* executable = NULL; + iree_hal_metal_executable_t* executable = NULL; IREE_RETURN_IF_ERROR( - iree_hal_metal_kernel_library_flatbuffer_verify(executable_params->executable_data)); + iree_hal_metal_executable_flatbuffer_verify(executable_params->executable_data)); iree_hal_metal_ExecutableDef_table_t executable_def = iree_hal_metal_ExecutableDef_as_root(executable_params->executable_data.data); @@ -276,7 +276,7 @@ iree_status_t iree_hal_metal_kernel_library_create( (char*)((char*)executable + sizeof(*executable) + entry_point_count * sizeof(executable->entry_points[0]))); if (iree_status_is_ok(status)) { - iree_hal_resource_initialize(&iree_hal_metal_kernel_library_vtable, &executable->resource); + iree_hal_resource_initialize(&iree_hal_metal_executable_vtable, &executable->resource); executable->host_allocator = host_allocator; executable->entry_point_count = entry_point_count; @@ -313,8 +313,8 @@ iree_status_t iree_hal_metal_kernel_library_create( } if (!iree_status_is_ok(status)) break; - status = iree_hal_metal_create_pipline_object(library, entry_point_view, source_code, device, - &function, &pso); + status = iree_hal_metal_create_pipeline_object(library, entry_point_view, source_code, device, + &function, &pso); if (!iree_status_is_ok(status)) break; // Package required parameters for kernel launches for each entry point. @@ -350,8 +350,8 @@ iree_status_t iree_hal_metal_kernel_library_create( return status; } -static void iree_hal_metal_kernel_library_destroy(iree_hal_executable_t* base_executable) { - iree_hal_metal_kernel_library_t* executable = iree_hal_metal_kernel_library_cast(base_executable); +static void iree_hal_metal_executable_destroy(iree_hal_executable_t* base_executable) { + iree_hal_metal_executable_t* executable = iree_hal_metal_executable_cast(base_executable); IREE_TRACE_ZONE_BEGIN(z0); for (iree_host_size_t i = 0; i < executable->entry_point_count; ++i) { @@ -366,11 +366,11 @@ static void iree_hal_metal_kernel_library_destroy(iree_hal_executable_t* base_ex IREE_TRACE_ZONE_END(z0); } -iree_status_t iree_hal_metal_kernel_library_entry_point_kernel_params( +iree_status_t iree_hal_metal_executable_entry_point_kernel_params( const iree_hal_executable_t* base_executable, int32_t entry_point, iree_hal_metal_kernel_params_t* out_params) { - const iree_hal_metal_kernel_library_t* executable = - iree_hal_metal_kernel_library_const_cast(base_executable); + const iree_hal_metal_executable_t* executable = + iree_hal_metal_executable_const_cast(base_executable); if (entry_point >= executable->entry_point_count) { return iree_make_status(IREE_STATUS_OUT_OF_RANGE, "invalid entry point ordinal %d", entry_point); @@ -379,6 +379,6 @@ iree_status_t iree_hal_metal_kernel_library_entry_point_kernel_params( return iree_ok_status(); } -static const iree_hal_executable_vtable_t iree_hal_metal_kernel_library_vtable = { - .destroy = iree_hal_metal_kernel_library_destroy, +static const iree_hal_executable_vtable_t iree_hal_metal_executable_vtable = { + .destroy = iree_hal_metal_executable_destroy, }; diff --git a/runtime/src/iree/hal/drivers/metal/nop_executable_cache.m b/runtime/src/iree/hal/drivers/metal/nop_executable_cache.m index 347ce7d72a7a..f57be2871156 100644 --- a/runtime/src/iree/hal/drivers/metal/nop_executable_cache.m +++ b/runtime/src/iree/hal/drivers/metal/nop_executable_cache.m @@ -11,7 +11,7 @@ #include "iree/base/api.h" #include "iree/base/tracing.h" -#include "iree/hal/drivers/metal/kernel_library.h" +#include "iree/hal/drivers/metal/executable.h" typedef struct iree_hal_metal_nop_executable_cache_t { // Abstract resource used for injecting reference counting and vtable; must be at offset 0. @@ -75,8 +75,8 @@ static iree_status_t iree_hal_metal_nop_executable_cache_prepare_executable( const iree_hal_executable_params_t* executable_params, iree_hal_executable_t** out_executable) { iree_hal_metal_nop_executable_cache_t* executable_cache = iree_hal_metal_nop_executable_cache_cast(base_executable_cache); - return iree_hal_metal_kernel_library_create(executable_cache->device, executable_params, - executable_cache->host_allocator, out_executable); + return iree_hal_metal_executable_create(executable_cache->device, executable_params, + executable_cache->host_allocator, out_executable); } static const iree_hal_executable_cache_vtable_t iree_hal_metal_nop_executable_cache_vtable = { From d938c719f2e665c4b141e66e288e906413dc8461 Mon Sep 17 00:00:00 2001 From: Ben Vanik Date: Tue, 13 Aug 2024 12:51:10 -0700 Subject: [PATCH 4/6] 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. --- build_tools/bazel/iree_flatcc.bzl | 6 +- .../bazel_to_cmake_converter.py | 4 +- build_tools/cmake/flatbuffer_c_library.cmake | 3 +- compiler/plugins/target/CUDA/BUILD.bazel | 2 + compiler/plugins/target/CUDA/CMakeLists.txt | 2 + compiler/plugins/target/CUDA/CUDATarget.cpp | 10 +- .../plugins/target/MetalSPIRV/BUILD.bazel | 2 + .../plugins/target/MetalSPIRV/CMakeLists.txt | 2 + .../target/MetalSPIRV/MetalSPIRVTarget.cpp | 7 ++ compiler/plugins/target/ROCM/BUILD.bazel | 2 + compiler/plugins/target/ROCM/CMakeLists.txt | 2 + compiler/plugins/target/ROCM/ROCMTarget.cpp | 44 +++------ .../plugins/target/VulkanSPIRV/BUILD.bazel | 2 + .../plugins/target/VulkanSPIRV/CMakeLists.txt | 2 + .../target/VulkanSPIRV/VulkanSPIRVTarget.cpp | 39 +++----- .../plugins/target/WebGPUSPIRV/CMakeLists.txt | 2 + .../target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp | 6 ++ .../compiler/Dialect/HAL/Utils/BUILD.bazel | 16 +++ .../compiler/Dialect/HAL/Utils/CMakeLists.txt | 15 +++ .../HAL/Utils/ExecutableDebugInfoUtils.cpp | 99 +++++++++++++++++++ .../HAL/Utils/ExecutableDebugInfoUtils.h | 43 ++++++++ experimental/webgpu/BUILD.bazel | 2 + experimental/webgpu/executable.c | 7 ++ runtime/src/iree/hal/drivers/cuda/BUILD.bazel | 2 + .../src/iree/hal/drivers/cuda/CMakeLists.txt | 2 + .../iree/hal/drivers/cuda/native_executable.c | 20 +++- .../src/iree/hal/drivers/hip/CMakeLists.txt | 2 + .../iree/hal/drivers/hip/native_executable.c | 19 ++-- .../src/iree/hal/drivers/metal/CMakeLists.txt | 2 + .../src/iree/hal/drivers/metal/executable.m | 9 ++ .../src/iree/hal/drivers/vulkan/BUILD.bazel | 2 + .../iree/hal/drivers/vulkan/CMakeLists.txt | 2 + .../hal/drivers/vulkan/native_executable.cc | 57 +++++------ runtime/src/iree/hal/utils/BUILD.bazel | 31 ++++-- runtime/src/iree/hal/utils/CMakeLists.txt | 38 ++++--- .../iree/hal/utils/executable_debug_info.c | 74 ++++++++++++++ .../iree/hal/utils/executable_debug_info.h | 36 +++++++ runtime/src/iree/schemas/BUILD.bazel | 18 +++- runtime/src/iree/schemas/CMakeLists.txt | 27 ++++- .../src/iree/schemas/cuda_executable_def.fbs | 15 ++- .../iree/schemas/executable_debug_info.fbs | 44 +++++++++ .../src/iree/schemas/metal_executable_def.fbs | 5 + .../src/iree/schemas/rocm_executable_def.fbs | 32 +----- .../src/iree/schemas/spirv_executable_def.fbs | 31 +----- .../src/iree/schemas/wgsl_executable_def.fbs | 5 + 45 files changed, 596 insertions(+), 196 deletions(-) create mode 100644 compiler/src/iree/compiler/Dialect/HAL/Utils/ExecutableDebugInfoUtils.cpp create mode 100644 compiler/src/iree/compiler/Dialect/HAL/Utils/ExecutableDebugInfoUtils.h create mode 100644 runtime/src/iree/hal/utils/executable_debug_info.c create mode 100644 runtime/src/iree/hal/utils/executable_debug_info.h create mode 100644 runtime/src/iree/schemas/executable_debug_info.fbs diff --git a/build_tools/bazel/iree_flatcc.bzl b/build_tools/bazel/iree_flatcc.bzl index 80e66d67c0d4..de5ddb72e0fa 100644 --- a/build_tools/bazel/iree_flatcc.bzl +++ b/build_tools/bazel/iree_flatcc.bzl @@ -10,12 +10,14 @@ def iree_flatbuffer_c_library( 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 @@ def iree_flatbuffer_c_library( 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 9a1796b443fd..eb5d2b1ddd67 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 @@ def iree_bytecode_module( 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 fe0913cc1b3d..2016cdf75fb0 100644 --- a/build_tools/cmake/flatbuffer_c_library.cmake +++ b/build_tools/cmake/flatbuffer_c_library.cmake @@ -48,7 +48,7 @@ function(flatbuffer_c_library) cmake_parse_arguments(_RULE "PUBLIC;TESTONLY" "NAME" - "SRCS;FLATCC_ARGS" + "SRCS;FLATCC_ARGS;INCLUDES" ${ARGN} ) @@ -94,6 +94,7 @@ function(flatbuffer_c_library) 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 2d475bf0f45c..b694187f7325 100644 --- a/compiler/plugins/target/CUDA/BUILD.bazel +++ b/compiler/plugins/target/CUDA/BUILD.bazel @@ -33,11 +33,13 @@ iree_compiler_cc_library( "//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 214f78b5270f..70c6dc6b8a5b 100644 --- a/compiler/plugins/target/CUDA/CMakeLists.txt +++ b/compiler/plugins/target/CUDA/CMakeLists.txt @@ -57,10 +57,12 @@ iree_cc_library( 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 eb05e187fb03..5a41ffe67642 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" @@ -517,9 +518,13 @@ class CUDATargetBackend final : public TargetBackend { FlatbufferBuilder builder; iree_hal_cuda_ExecutableDef_start_as_root(builder); + // Attach embedded source file contents. + auto sourceFilesRef = createSourceFilesVec( + serOptions.debugLevel, variantOp.getSourcesAttr(), builder); + SmallVector entryPointNames; std::string ptxImage; - SmallVector sourceLocationRefs; + SmallVector sourceLocationRefs; if (variantOp.isExternal()) { if (!variantOp.getObjects().has_value()) { return variantOp.emitOpError() @@ -590,7 +595,7 @@ class CUDATargetBackend final : public TargetBackend { 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())); } } @@ -686,6 +691,7 @@ class CUDATargetBackend final : public TargetBackend { 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 ede556649d07..ae750cbce17c 100644 --- a/compiler/plugins/target/MetalSPIRV/BUILD.bazel +++ b/compiler/plugins/target/MetalSPIRV/BUILD.bazel @@ -31,8 +31,10 @@ iree_compiler_cc_library( "//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 678a37a1a4d7..3a7b0e6f8a51 100644 --- a/compiler/plugins/target/MetalSPIRV/CMakeLists.txt +++ b/compiler/plugins/target/MetalSPIRV/CMakeLists.txt @@ -41,8 +41,10 @@ iree_cc_library( 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 4fa2b03c1094..8372307973ee 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 @@ class MetalSPIRVTargetBackend : public TargetBackend { 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 @@ class MetalSPIRVTargetBackend : public TargetBackend { 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 75296b81eb0c..47711208c9de 100644 --- a/compiler/plugins/target/ROCM/BUILD.bazel +++ b/compiler/plugins/target/ROCM/BUILD.bazel @@ -35,9 +35,11 @@ iree_compiler_cc_library( "//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 b3e8fd53e6a9..ca749feeb8c0 100644 --- a/compiler/plugins/target/ROCM/CMakeLists.txt +++ b/compiler/plugins/target/ROCM/CMakeLists.txt @@ -60,9 +60,11 @@ iree_cc_library( 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 ff9a3694f037..0465ee2b19c0 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" @@ -575,26 +576,11 @@ class ROCMTargetBackend final : public TargetBackend { iree_hal_rocm_ExecutableDef_start_as_root(builder); // Attach embedded source file contents. - SmallVector sourceFileRefs; - if (auto sourcesAttr = variantOp.getSourcesAttr()) { - for (auto sourceAttr : llvm::reverse(sourcesAttr.getValue())) { - if (auto resourceAttr = dyn_cast_if_present( - 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 entryPointNames; - SmallVector sourceLocationRefs; + SmallVector sourceLocationRefs; entryPointNames.resize(exportOps.size()); for (auto exportOp : exportOps) { auto ordinalAttr = exportOp.getOrdinalAttr(); @@ -612,27 +598,28 @@ class ROCMTargetBackend final : public TargetBackend { // 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 stageLocationsRefs; + SmallVector stageLocationsRefs; if (serOptions.debugLevel >= 3) { for (auto exportOp : exportOps) { - SmallVector stageLocationRefs; + SmallVector stageLocationRefs; if (auto locsAttr = exportOp.getSourceLocsAttr()) { for (auto locAttr : locsAttr.getValue()) { if (auto loc = findFirstFileLoc(cast(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()))); } } } @@ -641,7 +628,7 @@ class ROCMTargetBackend final : public TargetBackend { // 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)); } } @@ -678,10 +665,7 @@ class ROCMTargetBackend final : public TargetBackend { 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 984bef9c92d9..8419ed7ed4c6 100644 --- a/compiler/plugins/target/VulkanSPIRV/BUILD.bazel +++ b/compiler/plugins/target/VulkanSPIRV/BUILD.bazel @@ -29,8 +29,10 @@ iree_compiler_cc_library( "//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 958e27742d0e..b55461702180 100644 --- a/compiler/plugins/target/VulkanSPIRV/CMakeLists.txt +++ b/compiler/plugins/target/VulkanSPIRV/CMakeLists.txt @@ -37,8 +37,10 @@ iree_cc_library( 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 55fbea52e499..8ac4affb2af9 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" @@ -179,23 +180,8 @@ class VulkanSPIRVTargetBackend : public TargetBackend { iree_hal_spirv_ExecutableDef_start_as_root(builder); // Attach embedded source file contents. - SmallVector sourceFileRefs; - if (auto sourcesAttr = variantOp.getSourcesAttr()) { - for (auto sourceAttr : llvm::reverse(sourcesAttr.getValue())) { - if (auto resourceAttr = dyn_cast_if_present( - 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 shaderModuleRefs; @@ -206,7 +192,7 @@ class VulkanSPIRVTargetBackend : public TargetBackend { SmallVector entryPointNames; SmallVector subgroupSizes; SmallVector shaderModuleIndices; - SmallVector sourceLocationRefs; + SmallVector sourceLocationRefs; entryPointNames.resize(ordinalCount); subgroupSizes.resize(ordinalCount); shaderModuleIndices.resize(ordinalCount); @@ -270,17 +256,17 @@ class VulkanSPIRVTargetBackend : public TargetBackend { // 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 stageLocationsRefs; + SmallVector stageLocationsRefs; if (options.debugLevel >= 3) { for (auto exportOp : exportOps) { - SmallVector stageLocationRefs; + SmallVector stageLocationRefs; if (auto locsAttr = exportOp.getSourceLocsAttr()) { for (auto locAttr : locsAttr.getValue()) { if (auto loc = @@ -288,9 +274,9 @@ class VulkanSPIRVTargetBackend : public TargetBackend { 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()))); } } @@ -300,7 +286,7 @@ class VulkanSPIRVTargetBackend : public TargetBackend { // 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)); } } @@ -334,10 +320,7 @@ class VulkanSPIRVTargetBackend : public TargetBackend { 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 caf4460d1d51..4b7ef4e49a30 100644 --- a/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt +++ b/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt @@ -52,8 +52,10 @@ iree_cc_library( 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 8fd7c531316d..61d996510ff3 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 @@ class WebGPUSPIRVTargetBackend : public TargetBackend { 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 @@ class WebGPUSPIRVTargetBackend : public TargetBackend { 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 77458152445a..2a77d86addc6 100644 --- a/compiler/src/iree/compiler/Dialect/HAL/Utils/BUILD.bazel +++ b/compiler/src/iree/compiler/Dialect/HAL/Utils/BUILD.bazel @@ -12,6 +12,22 @@ package( licenses = ["notice"], # Apache 2.0 ) +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 = [ diff --git a/compiler/src/iree/compiler/Dialect/HAL/Utils/CMakeLists.txt b/compiler/src/iree/compiler/Dialect/HAL/Utils/CMakeLists.txt index 696c5f33fe08..22e773273188 100644 --- a/compiler/src/iree/compiler/Dialect/HAL/Utils/CMakeLists.txt +++ b/compiler/src/iree/compiler/Dialect/HAL/Utils/CMakeLists.txt @@ -10,6 +10,21 @@ iree_add_all_subdirs() +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 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 000000000000..cef943ef7811 --- /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 sourceFileRefs; + for (auto sourceAttr : llvm::reverse(sourcesAttr.getValue())) { + if (auto resourceAttr = dyn_cast_if_present( + 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 +createExportDefs(int debugLevel, + ArrayRef exportOps, + FlatbufferBuilder &fbb) { + if (debugLevel < 1) { + // No debug information. + return {}; + } + + SmallVector 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 stageLocationRefs; + if (auto locsAttr = exportOp.getSourceLocsAttr()) { + for (auto locAttr : locsAttr.getValue()) { + if (auto loc = + findFirstFileLoc(cast(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 000000000000..0a6cd02900d6 --- /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 +createExportDefs(int debugLevel, + ArrayRef 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 4e802e6c1de9..2906580f6090 100644 --- a/experimental/webgpu/BUILD.bazel +++ b/experimental/webgpu/BUILD.bazel @@ -53,8 +53,10 @@ iree_runtime_cc_library( "//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 ff3822517c8d..1929e413e39e 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 @@ iree_status_t iree_hal_webgpu_executable_create( 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 89fbe0ae088e..a2aa95fca402 100644 --- a/runtime/src/iree/hal/drivers/cuda/BUILD.bazel +++ b/runtime/src/iree/hal/drivers/cuda/BUILD.bazel @@ -66,11 +66,13 @@ iree_runtime_cc_library( "//runtime/src/iree/hal", "//runtime/src/iree/hal/utils:collective_batch", "//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: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 ebf5386ad7f0..d6b289862e32 100644 --- a/runtime/src/iree/hal/drivers/cuda/CMakeLists.txt +++ b/runtime/src/iree/hal/drivers/cuda/CMakeLists.txt @@ -63,11 +63,13 @@ iree_cc_library( iree::hal iree::hal::utils::collective_batch 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::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 18ab1e62e592..c3b32e01ecd2 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 @@ iree_status_t iree_hal_cuda_native_executable_create( 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_status_t iree_hal_cuda_native_executable_create( 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 d20e777d14e8..1b4c2a0d1d30 100644 --- a/runtime/src/iree/hal/drivers/hip/CMakeLists.txt +++ b/runtime/src/iree/hal/drivers/hip/CMakeLists.txt @@ -68,11 +68,13 @@ iree_cc_library( 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::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::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 5feab821e06d..9e75e12bac44 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 @@ iree_status_t iree_hal_hip_native_executable_create( 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_status_t iree_hal_hip_native_executable_create( 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 85356f30ec91..c4a98356d426 100644 --- a/runtime/src/iree/hal/drivers/metal/CMakeLists.txt +++ b/runtime/src/iree/hal/drivers/metal/CMakeLists.txt @@ -42,9 +42,11 @@ iree_cc_library( 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 b30ea157f056..0416960fa7df 100644 --- a/runtime/src/iree/hal/drivers/metal/executable.m +++ b/runtime/src/iree/hal/drivers/metal/executable.m @@ -9,9 +9,12 @@ #include #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 @@ iree_status_t iree_hal_metal_executable_create( 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 ce5b68b4e3eb..68aef348faeb 100644 --- a/runtime/src/iree/hal/drivers/vulkan/BUILD.bazel +++ b/runtime/src/iree/hal/drivers/vulkan/BUILD.bazel @@ -79,10 +79,12 @@ iree_runtime_cc_library( "//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 76e376f91d15..b495ae62ac65 100644 --- a/runtime/src/iree/hal/drivers/vulkan/CMakeLists.txt +++ b/runtime/src/iree/hal/drivers/vulkan/CMakeLists.txt @@ -74,10 +74,12 @@ iree_cc_library( 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 ebfd00626532..b44efe8c0cd6 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 @@ typedef struct iree_hal_vulkan_entry_point_t { 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 @@ iree_status_t iree_hal_vulkan_native_executable_create( } } + // 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 @@ void iree_hal_vulkan_native_executable_entry_point_source_location( 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 395fc7f93858..ef2de8b373e7 100644 --- a/runtime/src/iree/hal/utils/BUILD.bazel +++ b/runtime/src/iree/hal/utils/BUILD.bazel @@ -25,16 +25,6 @@ iree_runtime_cc_library( ], ) -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"], @@ -58,6 +48,16 @@ iree_runtime_cc_library( ], ) +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"], @@ -70,6 +70,17 @@ iree_runtime_cc_library( ], ) +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"], diff --git a/runtime/src/iree/hal/utils/CMakeLists.txt b/runtime/src/iree/hal/utils/CMakeLists.txt index 78263704dd2d..12fd89c33769 100644 --- a/runtime/src/iree/hal/utils/CMakeLists.txt +++ b/runtime/src/iree/hal/utils/CMakeLists.txt @@ -27,42 +27,42 @@ iree_cc_library( iree_cc_library( NAME - debug_allocator + collective_batch HDRS - "debug_allocator.h" + "collective_batch.h" SRCS - "debug_allocator.c" + "collective_batch.c" DEPS + ::resource_set iree::base + iree::base::internal::arena iree::hal PUBLIC ) iree_cc_library( NAME - collective_batch + caching_allocator HDRS - "collective_batch.h" + "caching_allocator.h" SRCS - "collective_batch.c" + "caching_allocator.c" DEPS - ::resource_set iree::base - iree::base::internal::arena + iree::base::internal::synchronization iree::hal PUBLIC ) iree_cc_library( NAME - caching_allocator + debug_allocator HDRS - "caching_allocator.h" + "debug_allocator.h" SRCS - "caching_allocator.c" + "debug_allocator.c" DEPS iree::base - iree::base::internal::synchronization iree::hal PUBLIC ) @@ -82,6 +82,20 @@ iree_cc_library( PUBLIC ) +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 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 000000000000..0cd214995512 --- /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 000000000000..bae9961e7965 --- /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 294c7931f3e2..2f0959bbebf5 100644 --- a/runtime/src/iree/schemas/BUILD.bazel +++ b/runtime/src/iree/schemas/BUILD.bazel @@ -30,11 +30,12 @@ iree_flatbuffer_c_library( 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 @@ iree_flatbuffer_c_library( 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 776616eaacbf..cfbb8508a686 100644 --- a/runtime/src/iree/schemas/CMakeLists.txt +++ b/runtime/src/iree/schemas/CMakeLists.txt @@ -33,14 +33,16 @@ flatbuffer_c_library( "--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 @@ flatbuffer_c_library( "--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 @@ flatbuffer_c_library( "--builder" "--verifier" "--json" + INCLUDES + "executable_debug_info.fbs" PUBLIC ) @@ -85,6 +106,8 @@ flatbuffer_c_library( "--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 b6713d493d7a..0abc40ae550d 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 @@ struct BlockSize { 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 @@ table ExecutableDef { // 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 000000000000..c3bceaa292c7 --- /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 dc727818d639..fd0330769985 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 @@ table ExecutableDef { 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 f368e1f43fda..6781115d3fce 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 @@ struct BlockSize { 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 @@ table ExecutableDef { // 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 4eaea8fc426b..a5aa17ed9505 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 @@ table ShaderModuleDef { 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 @@ table ExecutableDef { // 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 79c821f3c58b..bba8f4c43096 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 @@ table ExecutableDef { // 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; From 81ad11804f20165101bdea80882c7a11cd8ed11e Mon Sep 17 00:00:00 2001 From: Ben Vanik Date: Tue, 13 Aug 2024 13:52:36 -0700 Subject: [PATCH 5/6] Renaming rocm executable -> hip executable. --- compiler/plugins/target/ROCM/BUILD.bazel | 2 +- compiler/plugins/target/ROCM/CMakeLists.txt | 2 +- compiler/plugins/target/ROCM/ROCMTarget.cpp | 30 ++++++------- experimental/webgpu/executable.c | 2 +- .../src/iree/hal/drivers/hip/CMakeLists.txt | 2 +- .../iree/hal/drivers/hip/native_executable.c | 42 +++++++++---------- runtime/src/iree/schemas/BUILD.bazel | 10 ++--- runtime/src/iree/schemas/CMakeLists.txt | 8 ++-- ...cutable_def.fbs => hip_executable_def.fbs} | 8 ++-- 9 files changed, 53 insertions(+), 53 deletions(-) rename runtime/src/iree/schemas/{rocm_executable_def.fbs => hip_executable_def.fbs} (93%) diff --git a/compiler/plugins/target/ROCM/BUILD.bazel b/compiler/plugins/target/ROCM/BUILD.bazel index 47711208c9de..7962cf8e6073 100644 --- a/compiler/plugins/target/ROCM/BUILD.bazel +++ b/compiler/plugins/target/ROCM/BUILD.bazel @@ -40,7 +40,7 @@ iree_compiler_cc_library( "//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", + "//runtime/src/iree/schemas:hip_executable_def_c_fbs", "@llvm-project//llvm:AMDGPUCodeGen", "@llvm-project//llvm:Analysis", "@llvm-project//llvm:BitWriter", diff --git a/compiler/plugins/target/ROCM/CMakeLists.txt b/compiler/plugins/target/ROCM/CMakeLists.txt index ca749feeb8c0..9430dca4fc16 100644 --- a/compiler/plugins/target/ROCM/CMakeLists.txt +++ b/compiler/plugins/target/ROCM/CMakeLists.txt @@ -65,7 +65,7 @@ iree_cc_library( iree::compiler::PluginAPI iree::compiler::Utils iree::schemas::executable_debug_info_c_fbs - iree::schemas::rocm_executable_def_c_fbs + iree::schemas::hip_executable_def_c_fbs PUBLIC ) diff --git a/compiler/plugins/target/ROCM/ROCMTarget.cpp b/compiler/plugins/target/ROCM/ROCMTarget.cpp index 0465ee2b19c0..0e100c4a644b 100644 --- a/compiler/plugins/target/ROCM/ROCMTarget.cpp +++ b/compiler/plugins/target/ROCM/ROCMTarget.cpp @@ -24,7 +24,7 @@ #include "iree/compiler/Utils/FlatbufferUtils.h" #include "iree/compiler/Utils/ModuleUtils.h" #include "iree/compiler/Utils/ToolUtils.h" -#include "iree/schemas/rocm_executable_def_builder.h" +#include "iree/schemas/hip_executable_def_builder.h" #include "llvm/ADT/StringExtras.h" #include "llvm/ADT/StringRef.h" #include "llvm/Analysis/TargetTransformInfo.h" @@ -573,7 +573,7 @@ class ROCMTargetBackend final : public TargetBackend { } iree_compiler::FlatbufferBuilder builder; - iree_hal_rocm_ExecutableDef_start_as_root(builder); + iree_hal_hip_ExecutableDef_start_as_root(builder); // Attach embedded source file contents. auto sourceFilesRef = createSourceFilesVec( @@ -638,35 +638,35 @@ class ROCMTargetBackend final : public TargetBackend { targetHSACO.size()); auto entryPointsRef = builder.createStringVec(entryPointNames); - iree_hal_rocm_BlockSize_vec_start(builder); + iree_hal_hip_BlockSize_vec_start(builder); auto blockSizes = workgroupSizes.begin(); for (int i = 0, e = entryPointNames.size(); i < e; ++i) { - iree_hal_rocm_BlockSize_vec_push_create( + iree_hal_hip_BlockSize_vec_push_create( builder, (*blockSizes)[0], (*blockSizes)[1], (*blockSizes)[2]); ++blockSizes; } auto workgroupLocalMemoriesRef = builder.createInt32Vec(workgroupLocalMemories); - auto blockSizesRef = iree_hal_rocm_BlockSize_vec_end(builder); - iree_hal_rocm_ExecutableDef_entry_points_add(builder, entryPointsRef); - iree_hal_rocm_ExecutableDef_block_sizes_add(builder, blockSizesRef); - iree_hal_rocm_ExecutableDef_shared_memory_sizes_add( + auto blockSizesRef = iree_hal_hip_BlockSize_vec_end(builder); + iree_hal_hip_ExecutableDef_entry_points_add(builder, entryPointsRef); + iree_hal_hip_ExecutableDef_block_sizes_add(builder, blockSizesRef); + iree_hal_hip_ExecutableDef_shared_memory_sizes_add( builder, workgroupLocalMemoriesRef); - iree_hal_rocm_ExecutableDef_hsaco_image_add(builder, hsacoRef); + iree_hal_hip_ExecutableDef_hsaco_image_add(builder, hsacoRef); if (!sourceLocationRefs.empty()) { auto sourceLocationsRef = builder.createOffsetVecDestructive(sourceLocationRefs); - iree_hal_rocm_ExecutableDef_source_locations_add(builder, - sourceLocationsRef); + iree_hal_hip_ExecutableDef_source_locations_add(builder, + sourceLocationsRef); } if (!stageLocationsRefs.empty()) { auto stageLocationsRef = builder.createOffsetVecDestructive(stageLocationsRefs); - iree_hal_rocm_ExecutableDef_stage_locations_add(builder, - stageLocationsRef); + iree_hal_hip_ExecutableDef_stage_locations_add(builder, + stageLocationsRef); } - iree_hal_rocm_ExecutableDef_source_files_add(builder, sourceFilesRef); - iree_hal_rocm_ExecutableDef_end_as_root(builder); + iree_hal_hip_ExecutableDef_source_files_add(builder, sourceFilesRef); + iree_hal_hip_ExecutableDef_end_as_root(builder); // Add the binary data to the target executable. executableBuilder.create( diff --git a/experimental/webgpu/executable.c b/experimental/webgpu/executable.c index 1929e413e39e..0055e3bfe43e 100644 --- a/experimental/webgpu/executable.c +++ b/experimental/webgpu/executable.c @@ -273,7 +273,7 @@ iree_status_t iree_hal_webgpu_executable_create( // Publish any embedded source files to the tracing infrastructure. iree_hal_debug_publish_source_files( - iree_hal_rocm_ExecutableDef_source_files_get(executable_def)); + iree_hal_hip_ExecutableDef_source_files_get(executable_def)); // Create one pipeline per entry point. flatbuffers_uint32_vec_t entry_points_vec = diff --git a/runtime/src/iree/hal/drivers/hip/CMakeLists.txt b/runtime/src/iree/hal/drivers/hip/CMakeLists.txt index 1b4c2a0d1d30..f494cb8e73db 100644 --- a/runtime/src/iree/hal/drivers/hip/CMakeLists.txt +++ b/runtime/src/iree/hal/drivers/hip/CMakeLists.txt @@ -75,7 +75,7 @@ iree_cc_library( iree::hal::utils::resource_set iree::hal::utils::semaphore_base iree::schemas::executable_debug_info_c_fbs - iree::schemas::rocm_executable_def_c_fbs + iree::schemas::hip_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 9e75e12bac44..ce68229c890b 100644 --- a/runtime/src/iree/hal/drivers/hip/native_executable.c +++ b/runtime/src/iree/hal/drivers/hip/native_executable.c @@ -18,8 +18,8 @@ #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/rocm_executable_def_reader.h" -#include "iree/schemas/rocm_executable_def_verifier.h" +#include "iree/schemas/hip_executable_def_reader.h" +#include "iree/schemas/hip_executable_def_verifier.h" typedef struct iree_hal_hip_native_executable_t { // Abstract resource used for injecting reference counting and vtable; @@ -64,7 +64,7 @@ static iree_status_t iree_hal_hip_native_executable_flatbuffer_verify( // Run flatcc generated verification. This ensures all pointers are in-bounds // and that we can safely walk the file, but not that the actual contents of // the flatbuffer meet our expectations. - int verify_ret = iree_hal_rocm_ExecutableDef_verify_as_root( + int verify_ret = iree_hal_hip_ExecutableDef_verify_as_root( flatbuffer_data.data, flatbuffer_data.data_length); if (verify_ret != flatcc_verify_ok) { return iree_make_status(IREE_STATUS_INVALID_ARGUMENT, @@ -72,11 +72,11 @@ static iree_status_t iree_hal_hip_native_executable_flatbuffer_verify( flatcc_verify_error_string(verify_ret)); } - iree_hal_rocm_ExecutableDef_table_t executable_def = - iree_hal_rocm_ExecutableDef_as_root(flatbuffer_data.data); + iree_hal_hip_ExecutableDef_table_t executable_def = + iree_hal_hip_ExecutableDef_as_root(flatbuffer_data.data); flatbuffers_string_vec_t entry_points_vec = - iree_hal_rocm_ExecutableDef_entry_points_get(executable_def); + iree_hal_hip_ExecutableDef_entry_points_get(executable_def); size_t entry_point_count = flatbuffers_string_vec_len(entry_points_vec); if (entry_point_count == 0) { return iree_make_status(IREE_STATUS_INVALID_ARGUMENT, @@ -90,9 +90,9 @@ static iree_status_t iree_hal_hip_native_executable_flatbuffer_verify( } } - iree_hal_rocm_BlockSize_vec_t block_sizes_vec = - iree_hal_rocm_ExecutableDef_block_sizes_get(executable_def); - size_t block_size_count = iree_hal_rocm_BlockSize_vec_len(block_sizes_vec); + iree_hal_hip_BlockSize_vec_t block_sizes_vec = + iree_hal_hip_ExecutableDef_block_sizes_get(executable_def); + size_t block_size_count = iree_hal_hip_BlockSize_vec_len(block_sizes_vec); if (entry_point_count != block_size_count) { return iree_make_status( IREE_STATUS_INVALID_ARGUMENT, @@ -101,7 +101,7 @@ static iree_status_t iree_hal_hip_native_executable_flatbuffer_verify( } flatbuffers_uint32_vec_t shared_memory_sizes_vec = - iree_hal_rocm_ExecutableDef_shared_memory_sizes_get(executable_def); + iree_hal_hip_ExecutableDef_shared_memory_sizes_get(executable_def); size_t shared_memory_sizes_count = flatbuffers_string_vec_len(shared_memory_sizes_vec); if (entry_point_count != shared_memory_sizes_count) { @@ -112,7 +112,7 @@ static iree_status_t iree_hal_hip_native_executable_flatbuffer_verify( } flatbuffers_string_t hsaco_image = - iree_hal_rocm_ExecutableDef_hsaco_image_get(executable_def); + iree_hal_hip_ExecutableDef_hsaco_image_get(executable_def); if (flatbuffers_string_len(hsaco_image) == 0) { return iree_make_status(IREE_STATUS_INVALID_ARGUMENT, "no HSACO image present"); @@ -137,18 +137,18 @@ iree_status_t iree_hal_hip_native_executable_create( z0, iree_hal_hip_native_executable_flatbuffer_verify( executable_params->executable_data)); - iree_hal_rocm_ExecutableDef_table_t executable_def = - iree_hal_rocm_ExecutableDef_as_root( + iree_hal_hip_ExecutableDef_table_t executable_def = + iree_hal_hip_ExecutableDef_as_root( executable_params->executable_data.data); flatbuffers_string_vec_t entry_points_vec = - iree_hal_rocm_ExecutableDef_entry_points_get(executable_def); - iree_hal_rocm_BlockSize_vec_t block_sizes_vec = - iree_hal_rocm_ExecutableDef_block_sizes_get(executable_def); + iree_hal_hip_ExecutableDef_entry_points_get(executable_def); + iree_hal_hip_BlockSize_vec_t block_sizes_vec = + iree_hal_hip_ExecutableDef_block_sizes_get(executable_def); flatbuffers_uint32_vec_t shared_memory_sizes_vec = - iree_hal_rocm_ExecutableDef_shared_memory_sizes_get(executable_def); + iree_hal_hip_ExecutableDef_shared_memory_sizes_get(executable_def); flatbuffers_string_t hsaco_image = - iree_hal_rocm_ExecutableDef_hsaco_image_get(executable_def); + iree_hal_hip_ExecutableDef_hsaco_image_get(executable_def); iree_host_size_t entry_point_count = flatbuffers_string_vec_len(entry_points_vec); @@ -213,7 +213,7 @@ iree_status_t iree_hal_hip_native_executable_create( // Publish any embedded source files to the tracing infrastructure. iree_hal_debug_publish_source_files( - iree_hal_rocm_ExecutableDef_source_files_get(executable_def)); + iree_hal_hip_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 @@ -287,10 +287,10 @@ iree_status_t iree_hal_hip_native_executable_create( }); IREE_TRACE({ - if (iree_hal_rocm_ExecutableDef_source_locations_is_present( + if (iree_hal_hip_ExecutableDef_source_locations_is_present( executable_def)) { iree_hal_debug_FileLineLocDef_vec_t source_locs_vec = - iree_hal_rocm_ExecutableDef_source_locations_get(executable_def); + iree_hal_hip_ExecutableDef_source_locations_get(executable_def); iree_hal_debug_FileLineLocDef_table_t source_loc = iree_hal_debug_FileLineLocDef_vec_at(source_locs_vec, i); flatbuffers_string_t filename = diff --git a/runtime/src/iree/schemas/BUILD.bazel b/runtime/src/iree/schemas/BUILD.bazel index 2f0959bbebf5..324dc7dddc0c 100644 --- a/runtime/src/iree/schemas/BUILD.bazel +++ b/runtime/src/iree/schemas/BUILD.bazel @@ -40,15 +40,15 @@ iree_flatbuffer_c_library( ) iree_flatbuffer_c_library( - name = "metal_executable_def_c_fbs", - srcs = ["metal_executable_def.fbs"], + name = "hip_executable_def_c_fbs", + srcs = ["hip_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 = "metal_executable_def_c_fbs", + srcs = ["metal_executable_def.fbs"], flatcc_args = FLATCC_ARGS, includes = ["executable_debug_info.fbs"], ) @@ -73,8 +73,8 @@ iree_build_test( ":bytecode_module_def_c_fbs", ":cuda_executable_def_c_fbs", ":executable_debug_info_c_fbs", + ":hip_executable_def_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 cfbb8508a686..46647e4f6188 100644 --- a/runtime/src/iree/schemas/CMakeLists.txt +++ b/runtime/src/iree/schemas/CMakeLists.txt @@ -53,9 +53,9 @@ flatbuffer_c_library( flatbuffer_c_library( NAME - metal_executable_def_c_fbs + hip_executable_def_c_fbs SRCS - "metal_executable_def.fbs" + "hip_executable_def.fbs" FLATCC_ARGS "--reader" "--builder" @@ -68,9 +68,9 @@ flatbuffer_c_library( flatbuffer_c_library( NAME - rocm_executable_def_c_fbs + metal_executable_def_c_fbs SRCS - "rocm_executable_def.fbs" + "metal_executable_def.fbs" FLATCC_ARGS "--reader" "--builder" diff --git a/runtime/src/iree/schemas/rocm_executable_def.fbs b/runtime/src/iree/schemas/hip_executable_def.fbs similarity index 93% rename from runtime/src/iree/schemas/rocm_executable_def.fbs rename to runtime/src/iree/schemas/hip_executable_def.fbs index 6781115d3fce..55dc527daf26 100644 --- a/runtime/src/iree/schemas/rocm_executable_def.fbs +++ b/runtime/src/iree/schemas/hip_executable_def.fbs @@ -6,11 +6,11 @@ include "iree/schemas/executable_debug_info.fbs"; -namespace iree.hal.rocm; +namespace iree.hal.hip; -// 'ROCM Executable'. -file_identifier "ROCM"; -file_extension "rocm"; +// 'HIP Executable'. +file_identifier "HIPE"; +file_extension "hipe"; // A struct for the kernel block size along each dimensions. struct BlockSize { From ab732ab596dd1c9a33a63a12e73e326d09c424b8 Mon Sep 17 00:00:00 2001 From: Ben Vanik Date: Tue, 13 Aug 2024 17:44:27 -0700 Subject: [PATCH 6/6] Renaming [spirv|wgsl]_executable_def to [vulkan|webgpu]. I don't expect us to have the same flatbuffers for different HAL implementations and do expect us to have alternative file formats for the same HAL implementation. --- .../plugins/target/VulkanSPIRV/BUILD.bazel | 2 +- .../plugins/target/VulkanSPIRV/CMakeLists.txt | 2 +- .../target/VulkanSPIRV/VulkanSPIRVTarget.cpp | 44 ++++++------- .../plugins/target/WebGPUSPIRV/CMakeLists.txt | 2 +- .../target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp | 20 +++--- experimental/webgpu/BUILD.bazel | 2 +- experimental/webgpu/CMakeLists.txt | 2 +- experimental/webgpu/builtins.c | 4 +- experimental/webgpu/executable.c | 46 ++++++------- experimental/webgpu/shaders/BUILD.bazel | 2 +- experimental/webgpu/shaders/CMakeLists.txt | 2 +- .../src/iree/hal/drivers/vulkan/BUILD.bazel | 2 +- .../iree/hal/drivers/vulkan/CMakeLists.txt | 2 +- .../hal/drivers/vulkan/native_executable.cc | 66 +++++++++---------- runtime/src/iree/schemas/BUILD.bazel | 12 ++-- runtime/src/iree/schemas/CMakeLists.txt | 8 +-- .../src/iree/schemas/cuda_executable_def.fbs | 6 +- .../src/iree/schemas/hip_executable_def.fbs | 6 +- .../src/iree/schemas/metal_executable_def.fbs | 6 +- ...able_def.fbs => vulkan_executable_def.fbs} | 10 +-- ...able_def.fbs => webgpu_executable_def.fbs} | 10 +-- 21 files changed, 129 insertions(+), 127 deletions(-) rename runtime/src/iree/schemas/{spirv_executable_def.fbs => vulkan_executable_def.fbs} (93%) rename runtime/src/iree/schemas/{wgsl_executable_def.fbs => webgpu_executable_def.fbs} (88%) diff --git a/compiler/plugins/target/VulkanSPIRV/BUILD.bazel b/compiler/plugins/target/VulkanSPIRV/BUILD.bazel index 8419ed7ed4c6..fb53170d9f5c 100644 --- a/compiler/plugins/target/VulkanSPIRV/BUILD.bazel +++ b/compiler/plugins/target/VulkanSPIRV/BUILD.bazel @@ -33,7 +33,7 @@ iree_compiler_cc_library( "//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", + "//runtime/src/iree/schemas:vulkan_executable_def_c_fbs", "@llvm-project//llvm:Support", "@llvm-project//mlir:AsmParser", "@llvm-project//mlir:GPUDialect", diff --git a/compiler/plugins/target/VulkanSPIRV/CMakeLists.txt b/compiler/plugins/target/VulkanSPIRV/CMakeLists.txt index b55461702180..3ef8e7518711 100644 --- a/compiler/plugins/target/VulkanSPIRV/CMakeLists.txt +++ b/compiler/plugins/target/VulkanSPIRV/CMakeLists.txt @@ -41,7 +41,7 @@ iree_cc_library( iree::compiler::PluginAPI iree::compiler::Utils iree::schemas::executable_debug_info_c_fbs - iree::schemas::spirv_executable_def_c_fbs + iree::schemas::vulkan_executable_def_c_fbs PUBLIC ) diff --git a/compiler/plugins/target/VulkanSPIRV/VulkanSPIRVTarget.cpp b/compiler/plugins/target/VulkanSPIRV/VulkanSPIRVTarget.cpp index 8ac4affb2af9..b22b868fbff7 100644 --- a/compiler/plugins/target/VulkanSPIRV/VulkanSPIRVTarget.cpp +++ b/compiler/plugins/target/VulkanSPIRV/VulkanSPIRVTarget.cpp @@ -12,7 +12,7 @@ #include "iree/compiler/PluginAPI/Client.h" #include "iree/compiler/Utils/FlatbufferUtils.h" #include "iree/compiler/Utils/ModuleUtils.h" -#include "iree/schemas/spirv_executable_def_builder.h" +#include "iree/schemas/vulkan_executable_def_builder.h" #include "llvm/ADT/STLExtras.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/raw_ostream.h" @@ -177,14 +177,14 @@ class VulkanSPIRVTargetBackend : public TargetBackend { uint64_t ordinalCount = entryPointOrdinals.size(); FlatbufferBuilder builder; - iree_hal_spirv_ExecutableDef_start_as_root(builder); + iree_hal_vulkan_ExecutableDef_start_as_root(builder); // Attach embedded source file contents. auto sourceFilesRef = createSourceFilesVec( options.debugLevel, variantOp.getSourcesAttr(), builder); // The list of shader modules. - SmallVector shaderModuleRefs; + SmallVector shaderModuleRefs; // Per entry-point data. // Note that the following vectors should all be of the same size and @@ -232,7 +232,7 @@ class VulkanSPIRVTargetBackend : public TargetBackend { spvBinary.size()); shaderModuleIndices[ordinal] = shaderModuleRefs.size(); shaderModuleRefs.push_back( - iree_hal_spirv_ShaderModuleDef_create(builder, spvCodeRef)); + iree_hal_vulkan_ShaderModuleDef_create(builder, spvCodeRef)); // The IREE runtime uses ordinals instead of names. We need to attach the // entry point name for VkShaderModuleCreateInfo. @@ -298,31 +298,31 @@ class VulkanSPIRVTargetBackend : public TargetBackend { hasAnySubgroupSizes ? builder.createInt32Vec(subgroupSizes) : 0; flatbuffers_int32_vec_ref_t shaderModuleIndicesRef = builder.createInt32Vec(shaderModuleIndices); - iree_hal_spirv_ExecutableDef_entry_points_add(builder, entryPointsRef); + iree_hal_vulkan_ExecutableDef_entry_points_add(builder, entryPointsRef); if (subgroupSizesRef) { - iree_hal_spirv_ExecutableDef_subgroup_sizes_add(builder, - subgroupSizesRef); + iree_hal_vulkan_ExecutableDef_subgroup_sizes_add(builder, + subgroupSizesRef); } - iree_hal_spirv_ExecutableDef_shader_module_indices_add( + iree_hal_vulkan_ExecutableDef_shader_module_indices_add( builder, shaderModuleIndicesRef); auto shaderModulesRef = builder.createOffsetVecDestructive(shaderModuleRefs); - iree_hal_spirv_ExecutableDef_shader_modules_add(builder, shaderModulesRef); + iree_hal_vulkan_ExecutableDef_shader_modules_add(builder, shaderModulesRef); if (!sourceLocationRefs.empty()) { auto sourceLocationsRef = builder.createOffsetVecDestructive(sourceLocationRefs); - iree_hal_spirv_ExecutableDef_source_locations_add(builder, - sourceLocationsRef); + iree_hal_vulkan_ExecutableDef_source_locations_add(builder, + sourceLocationsRef); } if (!stageLocationsRefs.empty()) { auto stageLocationsRef = builder.createOffsetVecDestructive(stageLocationsRefs); - iree_hal_spirv_ExecutableDef_stage_locations_add(builder, - stageLocationsRef); + iree_hal_vulkan_ExecutableDef_stage_locations_add(builder, + stageLocationsRef); } - iree_hal_spirv_ExecutableDef_source_files_add(builder, sourceFilesRef); + iree_hal_vulkan_ExecutableDef_source_files_add(builder, sourceFilesRef); - iree_hal_spirv_ExecutableDef_end_as_root(builder); + iree_hal_vulkan_ExecutableDef_end_as_root(builder); // Add the binary data to the target executable. auto binaryOp = executableBuilder.create( @@ -375,26 +375,26 @@ class VulkanSPIRVTargetBackend : public TargetBackend { } FlatbufferBuilder builder; - iree_hal_spirv_ExecutableDef_start_as_root(builder); + iree_hal_vulkan_ExecutableDef_start_as_root(builder); auto spvCodeRef = flatbuffers_uint32_vec_create( builder, reinterpret_cast(spvBinary.data()), spvBinary.size() / sizeof(uint32_t)); - SmallVector shaderModuleRefs; + SmallVector shaderModuleRefs; shaderModuleRefs.push_back( - iree_hal_spirv_ShaderModuleDef_create(builder, spvCodeRef)); + iree_hal_vulkan_ShaderModuleDef_create(builder, spvCodeRef)); // Add top-level executable fields following their order of definition. auto entryPointsRef = builder.createStringVec(entryPointNames); auto shaderModuleIndicesRef = builder.createInt32Vec(shaderModuleIndices); - iree_hal_spirv_ExecutableDef_entry_points_add(builder, entryPointsRef); - iree_hal_spirv_ExecutableDef_shader_module_indices_add( + iree_hal_vulkan_ExecutableDef_entry_points_add(builder, entryPointsRef); + iree_hal_vulkan_ExecutableDef_shader_module_indices_add( builder, shaderModuleIndicesRef); auto shaderModulesRef = builder.createOffsetVecDestructive(shaderModuleRefs); - iree_hal_spirv_ExecutableDef_shader_modules_add(builder, shaderModulesRef); + iree_hal_vulkan_ExecutableDef_shader_modules_add(builder, shaderModulesRef); - iree_hal_spirv_ExecutableDef_end_as_root(builder); + iree_hal_vulkan_ExecutableDef_end_as_root(builder); // Add the binary data to the target executable. auto binaryOp = executableBuilder.create( diff --git a/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt b/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt index 4b7ef4e49a30..64b6b4cbe309 100644 --- a/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt +++ b/compiler/plugins/target/WebGPUSPIRV/CMakeLists.txt @@ -56,7 +56,7 @@ iree_cc_library( iree::compiler::PluginAPI iree::compiler::Utils iree::schemas::executable_debug_info_c_fbs - iree::schemas::wgsl_executable_def_c_fbs + iree::schemas::webgpu_executable_def_c_fbs libtint PUBLIC ) diff --git a/compiler/plugins/target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp b/compiler/plugins/target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp index 61d996510ff3..9323465060ec 100644 --- a/compiler/plugins/target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp +++ b/compiler/plugins/target/WebGPUSPIRV/WebGPUSPIRVTarget.cpp @@ -14,7 +14,7 @@ #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" +#include "iree/schemas/webgpu_executable_def_builder.h" #include "llvm/Support/CommandLine.h" #include "llvm/Support/FormatVariadic.h" #include "mlir/Dialect/Func/IR/FuncOps.h" @@ -237,28 +237,28 @@ class WebGPUSPIRVTargetBackend : public TargetBackend { // Pack the WGSL and metadata into a FlatBuffer. FlatbufferBuilder builder; - iree_hal_wgsl_ExecutableDef_start_as_root(builder); + iree_hal_webgpu_ExecutableDef_start_as_root(builder); // Attach embedded source file contents. auto sourceFilesRef = createSourceFilesVec( serOptions.debugLevel, variantOp.getSourcesAttr(), builder); - iree_hal_wgsl_ShaderModuleDef_start(builder); + iree_hal_webgpu_ShaderModuleDef_start(builder); auto wgslRef = builder.createString(wgsl.value()); - iree_hal_wgsl_ShaderModuleDef_code_add(builder, wgslRef); + iree_hal_webgpu_ShaderModuleDef_wgsl_source_add(builder, wgslRef); // TODO(scotttodd): populate source map - auto shaderModuleRef = iree_hal_wgsl_ShaderModuleDef_end(builder); + auto shaderModuleRef = iree_hal_webgpu_ShaderModuleDef_end(builder); - auto shaderModulesVec = iree_hal_wgsl_ShaderModuleDef_vec_create( + auto shaderModulesVec = iree_hal_webgpu_ShaderModuleDef_vec_create( builder, &shaderModuleRef, /*len=*/1); - iree_hal_wgsl_ExecutableDef_shader_modules_add(builder, shaderModulesVec); + iree_hal_webgpu_ExecutableDef_shader_modules_add(builder, shaderModulesVec); 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_webgpu_ExecutableDef_entry_points_add(builder, entryPointsRef); + iree_hal_webgpu_ExecutableDef_source_files_add(builder, sourceFilesRef); - iree_hal_wgsl_ExecutableDef_end_as_root(builder); + iree_hal_webgpu_ExecutableDef_end_as_root(builder); // Add the binary data to the target executable. auto binaryOp = executableBuilder.create( diff --git a/experimental/webgpu/BUILD.bazel b/experimental/webgpu/BUILD.bazel index 2906580f6090..c7cec08b3070 100644 --- a/experimental/webgpu/BUILD.bazel +++ b/experimental/webgpu/BUILD.bazel @@ -57,7 +57,7 @@ iree_runtime_cc_library( "//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", + "//runtime/src/iree/schemas:webgpu_executable_def_c_fbs", "@webgpu_headers", ], ) diff --git a/experimental/webgpu/CMakeLists.txt b/experimental/webgpu/CMakeLists.txt index 967e55b0e1e4..fa71067dfc69 100644 --- a/experimental/webgpu/CMakeLists.txt +++ b/experimental/webgpu/CMakeLists.txt @@ -50,7 +50,7 @@ iree_cc_library( iree::experimental::webgpu::shaders iree::hal::utils::file_transfer iree::hal::utils::memory_file - iree::schemas::wgsl_executable_def_c_fbs + iree::schemas::webgpu_executable_def_c_fbs PUBLIC ) diff --git a/experimental/webgpu/builtins.c b/experimental/webgpu/builtins.c index cfeced7a6a8d..de0d306587b5 100644 --- a/experimental/webgpu/builtins.c +++ b/experimental/webgpu/builtins.c @@ -10,8 +10,8 @@ #include "iree/base/api.h" static const char* iree_hal_webgpu_builtins_find_code(const char* file_name) { - const iree_file_toc_t* files = iree_hal_wgsl_builtin_shaders_create(); - for (size_t i = 0; i < iree_hal_wgsl_builtin_shaders_size(); ++i) { + const iree_file_toc_t* files = iree_hal_webgpu_builtin_shaders_create(); + for (size_t i = 0; i < iree_hal_webgpu_builtin_shaders_size(); ++i) { if (strcmp(file_name, files[i].name) == 0) { return files[i].data; } diff --git a/experimental/webgpu/executable.c b/experimental/webgpu/executable.c index 0055e3bfe43e..9191f5b73756 100644 --- a/experimental/webgpu/executable.c +++ b/experimental/webgpu/executable.c @@ -16,8 +16,8 @@ #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" +#include "iree/schemas/webgpu_executable_def_reader.h" +#include "iree/schemas/webgpu_executable_def_verifier.h" typedef struct iree_hal_webgpu_executable_t { iree_hal_resource_t resource; @@ -49,7 +49,7 @@ static iree_status_t iree_hal_webgpu_executable_flatbuffer_verify( // Run flatcc generated verification. This ensures all pointers are in-bounds // and that we can safely walk the file, but not that the actual contents of // the flatbuffer meet our expectations. - int verify_ret = iree_hal_wgsl_ExecutableDef_verify_as_root( + int verify_ret = iree_hal_webgpu_ExecutableDef_verify_as_root( flatbuffer_data.data, flatbuffer_data.data_length); if (verify_ret != flatcc_verify_ok) { return iree_make_status(IREE_STATUS_INVALID_ARGUMENT, @@ -57,18 +57,18 @@ static iree_status_t iree_hal_webgpu_executable_flatbuffer_verify( flatcc_verify_error_string(verify_ret)); } - iree_hal_wgsl_ExecutableDef_table_t executable_def = - iree_hal_wgsl_ExecutableDef_as_root(flatbuffer_data.data); + iree_hal_webgpu_ExecutableDef_table_t executable_def = + iree_hal_webgpu_ExecutableDef_as_root(flatbuffer_data.data); - iree_hal_wgsl_ShaderModuleDef_vec_t shader_modules_vec = - iree_hal_wgsl_ExecutableDef_shader_modules_get(executable_def); + iree_hal_webgpu_ShaderModuleDef_vec_t shader_modules_vec = + iree_hal_webgpu_ExecutableDef_shader_modules_get(executable_def); size_t shader_module_count = - iree_hal_wgsl_ShaderModuleDef_vec_len(shader_modules_vec); + iree_hal_webgpu_ShaderModuleDef_vec_len(shader_modules_vec); for (size_t i = 0; i < shader_module_count; ++i) { - iree_hal_wgsl_ShaderModuleDef_table_t shader_module_def = - iree_hal_wgsl_ShaderModuleDef_vec_at(shader_modules_vec, i); - if (flatbuffers_string_len( - iree_hal_wgsl_ShaderModuleDef_code_get(shader_module_def)) == 0) { + iree_hal_webgpu_ShaderModuleDef_table_t shader_module_def = + iree_hal_webgpu_ShaderModuleDef_vec_at(shader_modules_vec, i); + if (flatbuffers_string_len(iree_hal_webgpu_ShaderModuleDef_wgsl_source_get( + shader_module_def)) == 0) { return iree_make_status(IREE_STATUS_INVALID_ARGUMENT, "shader module %zu WGSL code is missing/empty", i); @@ -76,7 +76,7 @@ static iree_status_t iree_hal_webgpu_executable_flatbuffer_verify( } flatbuffers_uint32_vec_t entry_points_vec = - iree_hal_wgsl_ExecutableDef_entry_points_get(executable_def); + iree_hal_webgpu_ExecutableDef_entry_points_get(executable_def); size_t entry_point_count = flatbuffers_uint32_vec_len(entry_points_vec); if (entry_point_count != expected_entry_point_count) { return iree_make_status(IREE_STATUS_FAILED_PRECONDITION, @@ -99,14 +99,16 @@ static iree_status_t iree_hal_webgpu_executable_flatbuffer_verify( } static iree_status_t iree_hal_webgpu_create_wgsl_shader_module( - WGPUDevice device, iree_hal_wgsl_ShaderModuleDef_table_t shader_module_def, + WGPUDevice device, + iree_hal_webgpu_ShaderModuleDef_table_t shader_module_def, WGPUShaderModule* out_shader_module) { IREE_ASSERT_ARGUMENT(shader_module_def); IREE_ASSERT_ARGUMENT(out_shader_module); *out_shader_module = NULL; IREE_TRACE_ZONE_BEGIN(z0); - const char* code = iree_hal_wgsl_ShaderModuleDef_code_get(shader_module_def); + const char* code = + iree_hal_webgpu_ShaderModuleDef_wgsl_source_get(shader_module_def); const WGPUShaderModuleWGSLDescriptor descriptor = { .chain = @@ -232,17 +234,17 @@ iree_status_t iree_hal_webgpu_executable_create( z0, iree_hal_webgpu_executable_flatbuffer_verify( executable_params->executable_data, executable_params->pipeline_layout_count)); - iree_hal_wgsl_ExecutableDef_table_t executable_def = - iree_hal_wgsl_ExecutableDef_as_root( + iree_hal_webgpu_ExecutableDef_table_t executable_def = + iree_hal_webgpu_ExecutableDef_as_root( executable_params->executable_data.data); // Create shader modules. This will be cheap on some implementations like // Metal that need pipeline information in order to be JIT'ed from WGSL while // on others it can be more expensive. - iree_hal_wgsl_ShaderModuleDef_vec_t shader_modules_vec = - iree_hal_wgsl_ExecutableDef_shader_modules_get(executable_def); + iree_hal_webgpu_ShaderModuleDef_vec_t shader_modules_vec = + iree_hal_webgpu_ExecutableDef_shader_modules_get(executable_def); size_t shader_module_count = - iree_hal_wgsl_ShaderModuleDef_vec_len(shader_modules_vec); + iree_hal_webgpu_ShaderModuleDef_vec_len(shader_modules_vec); iree_inline_array(WGPUShaderModule, shader_modules, shader_module_count, host_allocator); memset(iree_inline_array_data(shader_modules), 0, @@ -250,7 +252,7 @@ iree_status_t iree_hal_webgpu_executable_create( iree_status_t status = iree_ok_status(); for (size_t i = 0; i < shader_module_count; ++i) { status = iree_hal_webgpu_create_wgsl_shader_module( - device, iree_hal_wgsl_ShaderModuleDef_vec_at(shader_modules_vec, i), + device, iree_hal_webgpu_ShaderModuleDef_vec_at(shader_modules_vec, i), iree_inline_array_at(shader_modules, i)); if (!iree_status_is_ok(status)) break; } @@ -277,7 +279,7 @@ iree_status_t iree_hal_webgpu_executable_create( // Create one pipeline per entry point. flatbuffers_uint32_vec_t entry_points_vec = - iree_hal_wgsl_ExecutableDef_entry_points_get(executable_def); + iree_hal_webgpu_ExecutableDef_entry_points_get(executable_def); for (iree_host_size_t i = 0; i < executable->entry_point_count; i++) { uint32_t module_ordinal = flatbuffers_uint32_vec_at(entry_points_vec, i); status = iree_hal_webgpu_create_pipeline( diff --git a/experimental/webgpu/shaders/BUILD.bazel b/experimental/webgpu/shaders/BUILD.bazel index bc6077eeda44..f42036167f7a 100644 --- a/experimental/webgpu/shaders/BUILD.bazel +++ b/experimental/webgpu/shaders/BUILD.bazel @@ -20,5 +20,5 @@ iree_c_embed_data( c_file_output = "builtin_shaders.c", flatten = True, h_file_output = "builtin_shaders.h", - identifier = "iree_hal_wgsl_builtin_shaders", + identifier = "iree_hal_webgpu_builtin_shaders", ) diff --git a/experimental/webgpu/shaders/CMakeLists.txt b/experimental/webgpu/shaders/CMakeLists.txt index 78cbc8cdec10..04cc4577c8fd 100644 --- a/experimental/webgpu/shaders/CMakeLists.txt +++ b/experimental/webgpu/shaders/CMakeLists.txt @@ -20,7 +20,7 @@ iree_c_embed_data( H_FILE_OUTPUT "builtin_shaders.h" IDENTIFIER - "iree_hal_wgsl_builtin_shaders" + "iree_hal_webgpu_builtin_shaders" FLATTEN PUBLIC ) diff --git a/runtime/src/iree/hal/drivers/vulkan/BUILD.bazel b/runtime/src/iree/hal/drivers/vulkan/BUILD.bazel index 68aef348faeb..13d4916fed7d 100644 --- a/runtime/src/iree/hal/drivers/vulkan/BUILD.bazel +++ b/runtime/src/iree/hal/drivers/vulkan/BUILD.bazel @@ -85,7 +85,7 @@ iree_runtime_cc_library( "//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", + "//runtime/src/iree/schemas:vulkan_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 b495ae62ac65..f07d435a9fc5 100644 --- a/runtime/src/iree/hal/drivers/vulkan/CMakeLists.txt +++ b/runtime/src/iree/hal/drivers/vulkan/CMakeLists.txt @@ -80,7 +80,7 @@ iree_cc_library( iree::hal::utils::resource_set iree::hal::utils::semaphore_base iree::schemas::executable_debug_info_c_fbs - iree::schemas::spirv_executable_def_c_fbs + iree::schemas::vulkan_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 b44efe8c0cd6..a1bbc5a987ac 100644 --- a/runtime/src/iree/hal/drivers/vulkan/native_executable.cc +++ b/runtime/src/iree/hal/drivers/vulkan/native_executable.cc @@ -22,8 +22,8 @@ #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" +#include "iree/schemas/vulkan_executable_def_reader.h" +#include "iree/schemas/vulkan_executable_def_verifier.h" using namespace iree::hal::vulkan; @@ -66,7 +66,7 @@ static void iree_hal_vulkan_destroy_shader_module( static iree_status_t iree_hal_vulkan_create_pipelines( VkDeviceHandle* logical_device, VkPipelineCache pipeline_cache, const iree_hal_executable_params_t* executable_params, - iree_hal_spirv_ExecutableDef_table_t executable_def, + iree_hal_vulkan_ExecutableDef_table_t executable_def, VkShaderModule* shader_modules, iree_host_size_t pipeline_count, iree_hal_vulkan_entry_point_t* out_entry_points) { IREE_TRACE_SCOPE(); @@ -104,11 +104,11 @@ static iree_status_t iree_hal_vulkan_create_pipelines( } flatbuffers_string_vec_t entry_points_vec = - iree_hal_spirv_ExecutableDef_entry_points_get(executable_def); + iree_hal_vulkan_ExecutableDef_entry_points_get(executable_def); flatbuffers_uint32_vec_t shader_module_indices_vec = - iree_hal_spirv_ExecutableDef_shader_module_indices_get(executable_def); + iree_hal_vulkan_ExecutableDef_shader_module_indices_get(executable_def); flatbuffers_uint32_vec_t subgroup_sizes_vec = - iree_hal_spirv_ExecutableDef_subgroup_sizes_get(executable_def); + iree_hal_vulkan_ExecutableDef_subgroup_sizes_get(executable_def); for (iree_host_size_t entry_ordinal = 0; entry_ordinal < pipeline_count; ++entry_ordinal) { iree_hal_pipeline_layout_t* pipeline_layout = @@ -205,7 +205,7 @@ static void iree_hal_vulkan_destroy_pipeline(VkDeviceHandle* logical_device, // runtime. There are still some conditions we must be aware of (such as omitted // names on functions with internal linkage), however we shouldn't need to // bounds check anything within the FlatBuffer after this succeeds. -static iree_status_t iree_hal_spirv_executable_flatbuffer_verify( +static iree_status_t iree_hal_vulkan_executable_flatbuffer_verify( iree_const_byte_span_t flatbuffer_data, iree_host_size_t expected_entry_point_count) { if (!flatbuffer_data.data || flatbuffer_data.data_length < 16) { @@ -219,7 +219,7 @@ static iree_status_t iree_hal_spirv_executable_flatbuffer_verify( // Run flatcc generated verification. This ensures all pointers are in-bounds // and that we can safely walk the file, but not that the actual contents of // the FlatBuffer meet our expectations. - int verify_ret = iree_hal_spirv_ExecutableDef_verify_as_root( + int verify_ret = iree_hal_vulkan_ExecutableDef_verify_as_root( flatbuffer_data.data, flatbuffer_data.data_length); if (verify_ret != flatcc_verify_ok) { return iree_make_status(IREE_STATUS_INVALID_ARGUMENT, @@ -227,11 +227,11 @@ static iree_status_t iree_hal_spirv_executable_flatbuffer_verify( flatcc_verify_error_string(verify_ret)); } - iree_hal_spirv_ExecutableDef_table_t executable_def = - iree_hal_spirv_ExecutableDef_as_root(flatbuffer_data.data); + iree_hal_vulkan_ExecutableDef_table_t executable_def = + iree_hal_vulkan_ExecutableDef_as_root(flatbuffer_data.data); flatbuffers_string_vec_t entry_points_vec = - iree_hal_spirv_ExecutableDef_entry_points_get(executable_def); + iree_hal_vulkan_ExecutableDef_entry_points_get(executable_def); size_t entry_point_count = flatbuffers_string_vec_len(entry_points_vec); if (entry_point_count != expected_entry_point_count) { return iree_make_status(IREE_STATUS_FAILED_PRECONDITION, @@ -249,7 +249,7 @@ static iree_status_t iree_hal_spirv_executable_flatbuffer_verify( } flatbuffers_uint32_vec_t subgroup_sizes_vec = - iree_hal_spirv_ExecutableDef_subgroup_sizes_get(executable_def); + iree_hal_vulkan_ExecutableDef_subgroup_sizes_get(executable_def); if (subgroup_sizes_vec) { size_t subgroup_sizes_count = flatbuffers_vec_len(subgroup_sizes_vec); if (subgroup_sizes_count != expected_entry_point_count) { @@ -261,18 +261,18 @@ static iree_status_t iree_hal_spirv_executable_flatbuffer_verify( } } - iree_hal_spirv_ShaderModuleDef_vec_t shader_modules_vec = - iree_hal_spirv_ExecutableDef_shader_modules_get(executable_def); + iree_hal_vulkan_ShaderModuleDef_vec_t shader_modules_vec = + iree_hal_vulkan_ExecutableDef_shader_modules_get(executable_def); size_t shader_module_count = flatbuffers_vec_len(shader_modules_vec); if (shader_module_count == 0) { return iree_make_status(IREE_STATUS_FAILED_PRECONDITION, "executable provides no shader modules"); } for (size_t i = 0; i < shader_module_count; ++i) { - iree_hal_spirv_ShaderModuleDef_table_t shader_module = - iree_hal_spirv_ShaderModuleDef_vec_at(shader_modules_vec, i); + iree_hal_vulkan_ShaderModuleDef_table_t shader_module = + iree_hal_vulkan_ShaderModuleDef_vec_at(shader_modules_vec, i); size_t code_size = flatbuffers_uint32_vec_len( - iree_hal_spirv_ShaderModuleDef_code_get(shader_module)); + iree_hal_vulkan_ShaderModuleDef_spirv_code_get(shader_module)); if (code_size == 0) { return iree_make_status( IREE_STATUS_INVALID_ARGUMENT, @@ -281,7 +281,7 @@ static iree_status_t iree_hal_spirv_executable_flatbuffer_verify( } flatbuffers_uint32_vec_t shader_module_indices_vec = - iree_hal_spirv_ExecutableDef_shader_module_indices_get(executable_def); + iree_hal_vulkan_ExecutableDef_shader_module_indices_get(executable_def); size_t shader_module_index_count = flatbuffers_vec_len(shader_module_indices_vec); if (shader_module_index_count != expected_entry_point_count) { @@ -339,16 +339,16 @@ iree_status_t iree_hal_vulkan_native_executable_create( // Verify and fetch the executable FlatBuffer wrapper. IREE_RETURN_AND_END_ZONE_IF_ERROR( - z0, iree_hal_spirv_executable_flatbuffer_verify( + z0, iree_hal_vulkan_executable_flatbuffer_verify( executable_params->executable_data, executable_params->pipeline_layout_count)); - iree_hal_spirv_ExecutableDef_table_t executable_def = - iree_hal_spirv_ExecutableDef_as_root( + iree_hal_vulkan_ExecutableDef_table_t executable_def = + iree_hal_vulkan_ExecutableDef_as_root( executable_params->executable_data.data); // Allocate space for Vulkan shader module handles. - iree_hal_spirv_ShaderModuleDef_vec_t shader_modules_vec = - iree_hal_spirv_ExecutableDef_shader_modules_get(executable_def); + iree_hal_vulkan_ShaderModuleDef_vec_t shader_modules_vec = + iree_hal_vulkan_ExecutableDef_shader_modules_get(executable_def); size_t shader_module_count = flatbuffers_vec_len(shader_modules_vec); VkShaderModule* shader_modules = NULL; IREE_RETURN_AND_END_ZONE_IF_ERROR( @@ -360,10 +360,10 @@ iree_status_t iree_hal_vulkan_native_executable_create( // TODO: perform the shader module creation in multiple threaded manner. iree_status_t status = iree_ok_status(); for (size_t i = 0; i < shader_module_count; ++i) { - iree_hal_spirv_ShaderModuleDef_table_t shader_module = - iree_hal_spirv_ShaderModuleDef_vec_at(shader_modules_vec, i); + iree_hal_vulkan_ShaderModuleDef_table_t shader_module = + iree_hal_vulkan_ShaderModuleDef_vec_at(shader_modules_vec, i); flatbuffers_uint32_vec_t code_vec = - iree_hal_spirv_ShaderModuleDef_code_get(shader_module); + iree_hal_vulkan_ShaderModuleDef_spirv_code_get(shader_module); size_t code_size = flatbuffers_uint32_vec_len(code_vec) * sizeof(uint32_t); status = iree_hal_vulkan_create_shader_module( logical_device, iree_make_const_byte_span(code_vec, code_size), @@ -373,7 +373,7 @@ iree_status_t iree_hal_vulkan_native_executable_create( // Create pipelines for each entry point. flatbuffers_string_vec_t entry_points_vec = - iree_hal_spirv_ExecutableDef_entry_points_get(executable_def); + iree_hal_vulkan_ExecutableDef_entry_points_get(executable_def); iree_host_size_t entry_point_count = flatbuffers_string_vec_len(entry_points_vec); @@ -407,7 +407,7 @@ iree_status_t iree_hal_vulkan_native_executable_create( if (iree_status_is_ok(status)) { flatbuffers_string_vec_t entry_points_vec = - iree_hal_spirv_ExecutableDef_entry_points_get(executable_def); + iree_hal_vulkan_ExecutableDef_entry_points_get(executable_def); for (iree_host_size_t i = 0; i < entry_point_count; ++i) { flatbuffers_string_t name = flatbuffers_string_vec_at(entry_points_vec, i); @@ -420,24 +420,24 @@ iree_status_t iree_hal_vulkan_native_executable_create( // 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)); + iree_hal_vulkan_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( + if (iree_hal_vulkan_ExecutableDef_source_locations_is_present( executable_def)) { iree_hal_debug_FileLineLocDef_vec_t source_locations_vec = - iree_hal_spirv_ExecutableDef_source_locations_get(executable_def); + iree_hal_vulkan_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_debug_FileLineLocDef_vec_at(source_locations_vec, i); } } - if (iree_hal_spirv_ExecutableDef_stage_locations_is_present( + if (iree_hal_vulkan_ExecutableDef_stage_locations_is_present( executable_def)) { iree_hal_debug_StageLocationsDef_vec_t stage_locations_vec = - iree_hal_spirv_ExecutableDef_stage_locations_get(executable_def); + iree_hal_vulkan_ExecutableDef_stage_locations_get(executable_def); for (iree_host_size_t i = 0; i < entry_point_count; ++i) { iree_hal_debug_StageLocationsDef_table_t stage_locations = iree_hal_debug_StageLocationsDef_vec_at(stage_locations_vec, i); diff --git a/runtime/src/iree/schemas/BUILD.bazel b/runtime/src/iree/schemas/BUILD.bazel index 324dc7dddc0c..a8fbfcab8b12 100644 --- a/runtime/src/iree/schemas/BUILD.bazel +++ b/runtime/src/iree/schemas/BUILD.bazel @@ -54,15 +54,15 @@ iree_flatbuffer_c_library( ) iree_flatbuffer_c_library( - name = "spirv_executable_def_c_fbs", - srcs = ["spirv_executable_def.fbs"], + name = "vulkan_executable_def_c_fbs", + srcs = ["vulkan_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"], + name = "webgpu_executable_def_c_fbs", + srcs = ["webgpu_executable_def.fbs"], flatcc_args = FLATCC_ARGS, includes = ["executable_debug_info.fbs"], ) @@ -75,8 +75,8 @@ iree_build_test( ":executable_debug_info_c_fbs", ":hip_executable_def_c_fbs", ":metal_executable_def_c_fbs", - ":spirv_executable_def_c_fbs", - ":wgsl_executable_def_c_fbs", + ":vulkan_executable_def_c_fbs", + ":webgpu_executable_def_c_fbs", ], ) diff --git a/runtime/src/iree/schemas/CMakeLists.txt b/runtime/src/iree/schemas/CMakeLists.txt index 46647e4f6188..574b2cac4578 100644 --- a/runtime/src/iree/schemas/CMakeLists.txt +++ b/runtime/src/iree/schemas/CMakeLists.txt @@ -83,9 +83,9 @@ flatbuffer_c_library( flatbuffer_c_library( NAME - spirv_executable_def_c_fbs + vulkan_executable_def_c_fbs SRCS - "spirv_executable_def.fbs" + "vulkan_executable_def.fbs" FLATCC_ARGS "--reader" "--builder" @@ -98,9 +98,9 @@ flatbuffer_c_library( flatbuffer_c_library( NAME - wgsl_executable_def_c_fbs + webgpu_executable_def_c_fbs SRCS - "wgsl_executable_def.fbs" + "webgpu_executable_def.fbs" FLATCC_ARGS "--reader" "--builder" diff --git a/runtime/src/iree/schemas/cuda_executable_def.fbs b/runtime/src/iree/schemas/cuda_executable_def.fbs index 0abc40ae550d..0ba9c2552e38 100644 --- a/runtime/src/iree/schemas/cuda_executable_def.fbs +++ b/runtime/src/iree/schemas/cuda_executable_def.fbs @@ -8,9 +8,9 @@ include "iree/schemas/executable_debug_info.fbs"; namespace iree.hal.cuda; -// 'CUDA Executable'. -file_identifier "CUDA"; -file_extension "cuda"; +// 'CUDA v1 Executable'. +file_identifier "CDA1"; +file_extension "cda1"; // A struct for the kernel block size along each dimensions. struct BlockSize { diff --git a/runtime/src/iree/schemas/hip_executable_def.fbs b/runtime/src/iree/schemas/hip_executable_def.fbs index 55dc527daf26..2d5458b53078 100644 --- a/runtime/src/iree/schemas/hip_executable_def.fbs +++ b/runtime/src/iree/schemas/hip_executable_def.fbs @@ -8,9 +8,9 @@ include "iree/schemas/executable_debug_info.fbs"; namespace iree.hal.hip; -// 'HIP Executable'. -file_identifier "HIPE"; -file_extension "hipe"; +// 'HIP v1 Executable'. +file_identifier "HIP1"; +file_extension "hip1"; // A struct for the kernel block size along each dimensions. struct BlockSize { diff --git a/runtime/src/iree/schemas/metal_executable_def.fbs b/runtime/src/iree/schemas/metal_executable_def.fbs index fd0330769985..816123276903 100644 --- a/runtime/src/iree/schemas/metal_executable_def.fbs +++ b/runtime/src/iree/schemas/metal_executable_def.fbs @@ -8,9 +8,9 @@ include "iree/schemas/executable_debug_info.fbs"; namespace iree.hal.metal; -// 'Metal Executable'. -file_identifier "MTLE"; -file_extension "mtle"; +// 'Metal v1 Executable'. +file_identifier "MTL1"; +file_extension "mtl1"; // A struct for Metal threadgroup size along each dimension. struct ThreadgroupSize { diff --git a/runtime/src/iree/schemas/spirv_executable_def.fbs b/runtime/src/iree/schemas/vulkan_executable_def.fbs similarity index 93% rename from runtime/src/iree/schemas/spirv_executable_def.fbs rename to runtime/src/iree/schemas/vulkan_executable_def.fbs index a5aa17ed9505..3bd4e6d57c18 100644 --- a/runtime/src/iree/schemas/spirv_executable_def.fbs +++ b/runtime/src/iree/schemas/vulkan_executable_def.fbs @@ -6,15 +6,15 @@ include "iree/schemas/executable_debug_info.fbs"; -namespace iree.hal.spirv; +namespace iree.hal.vulkan; -// 'SPIR-V Executable'. -file_identifier "SPVE"; -file_extension "spve"; +// 'Vulkan v1 Executable'. +file_identifier "VKE1"; +file_extension "vk1"; table ShaderModuleDef { // SPIR-V code blob. - code:[uint32]; + spirv_code:[uint32]; } // A SPIR-V shader module and runtime pipeline layout description. diff --git a/runtime/src/iree/schemas/wgsl_executable_def.fbs b/runtime/src/iree/schemas/webgpu_executable_def.fbs similarity index 88% rename from runtime/src/iree/schemas/wgsl_executable_def.fbs rename to runtime/src/iree/schemas/webgpu_executable_def.fbs index bba8f4c43096..accdaa150a4c 100644 --- a/runtime/src/iree/schemas/wgsl_executable_def.fbs +++ b/runtime/src/iree/schemas/webgpu_executable_def.fbs @@ -6,18 +6,18 @@ include "iree/schemas/executable_debug_info.fbs"; -namespace iree.hal.wgsl; +namespace iree.hal.webgpu; -// 'WGSL Executable'. -file_identifier "WGSL"; -file_extension "wgsl"; +// 'WGSL v1 Executable'. +file_identifier "WGS1"; +file_extension "wgs1"; // Contents of one WGPUShaderModule, possibly with multiple entry points. // Entry points have the name "dN" where N is the executable-wide entry point // ordinal. table ShaderModuleDef { // WGSL source code. - code:string; + wgsl_source:string; // Optional `source-map-v3` format source map. source_map:string;