| // RUN: iree-opt --pass-pipeline='builtin.module(hal.executable(iree-hal-serialize-target-executables{target=rocm}))' \ |
| // RUN: --verify-diagnostics %s -o - |
| |
| // The final bitcode validation should error out on any external functions that |
| // remain in the final bitcode (post device bitcode linking). |
| |
| #executable_target_rocm_hsaco_fb = #hal.executable.target<"rocm", "rocm-hsaco-fb", |
| {iree.gpu.target = #iree_gpu.target<arch = "gfx942", features = "", |
| wgp = <compute = fp16, storage = b16, |
| subgroup = none, dot = none, mma = [], |
| subgroup_size_choices = [64], |
| max_workgroup_sizes = [1024, 1024, 1024], |
| max_thread_count_per_workgroup = 1024, |
| max_workgroup_memory_bytes = 65536, |
| max_workgroup_counts = [2147483647, 2147483647, 2147483647]>>, |
| ukernels = "none"}> |
| #pipeline_layout = #hal.pipeline.layout<bindings = [#hal.pipeline.binding<storage_buffer, Indirect>], |
| flags = Indirect> |
| builtin.module { |
| hal.executable public @test { |
| // expected-error @+2 {{found an unresolved external function 'external_func' in the final bitcode}} |
| // expected-error @+1 {{failed to serialize executable for target backend rocm}} |
| hal.executable.variant public @rocm_hsaco_fb target(#executable_target_rocm_hsaco_fb) { |
| hal.executable.export public @test ordinal(0) layout(#pipeline_layout) |
| attributes {subgroup_size = 64 : index, workgroup_size = [128 : index, 2 : index, 1 : index]} { |
| ^bb0(%arg0: !hal.device): |
| %c128 = arith.constant 128 : index |
| %c2 = arith.constant 2 : index |
| %c1 = arith.constant 1 : index |
| hal.return %c128, %c2, %c1 : index, index, index |
| } |
| builtin.module { |
| llvm.func @external_func() attributes {sym_visibility = "private"} |
| llvm.func @test() attributes { rocdl.kernel } { |
| llvm.call @external_func() : () -> () |
| llvm.return |
| } |
| } |
| } |
| } |
| } |
