// 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