| // Tests linking modules containing hal.executable.source ops. |
| // hal.executable.source is used for hand-authored extern dispatch kernels |
| // (e.g., pre-compiled GPU kernel objects). The linker must copy these |
| // alongside the functions that reference them via flow.dispatch. |
| |
| // RUN: iree-link \ |
| // RUN: --link-module=%p/iree-link-executable-source-module-d.mlir \ |
| // RUN: %s | FileCheck %s |
| |
| // External reference to module D's transform function. |
| util.func private @module_d.transform(%arg0: tensor<4xf32>) -> tensor<4xf32> |
| |
| // Main entry point that calls the linked function. |
| util.func public @main(%arg0: tensor<4xf32>) -> tensor<4xf32> { |
| %0 = util.call @module_d.transform(%arg0) : (tensor<4xf32>) -> tensor<4xf32> |
| util.return %0 : tensor<4xf32> |
| } |
| |
| // After linking: |
| // 1. The function body is filled in (not just a declaration). |
| // CHECK-DAG: util.func public @main |
| // CHECK-DAG: util.func private @module_d.transform |
| |
| // 2. The flow.dispatch reference in the linked function body is intact. |
| // CHECK-DAG: flow.dispatch @extern_kernel::@entry |
| |
| // 3. The hal.executable.source is copied with its body intact. |
| // CHECK-DAG: hal.executable.source private @extern_kernel |
| // CHECK-DAG: hal.executable.export public @entry ordinal(0) |
| // CHECK-DAG: workgroup_size = [64 : index, 1 : index, 1 : index] |
| |
| // ----- |
| |
| // Test that hal.executable.source internal symbols are NOT treated as |
| // external dependencies. The nested hal.executable.export should not |
| // appear as a linkable symbol. |
| |
| // RUN: iree-link --list-symbols \ |
| // RUN: %p/iree-link-executable-source-module-d.mlir \ |
| // RUN: | FileCheck %s --check-prefix=SYMBOLS |
| |
| // Only the public function should be listed, not the executable or its exports. |
| // SYMBOLS: Public symbols in |
| // SYMBOLS: @transform |
| // SYMBOLS-NOT: @extern_kernel |
| // SYMBOLS-NOT: @entry |
