[INFO] fetching crate svod-codegen 0.2.0...
[INFO] testing svod-codegen-0.2.0 against 1.100.0-beta.1+cargoflags=--release for beta-release-1.100-2
[INFO] extracting crate svod-codegen 0.2.0 into /workspace/builds/worker-7-tc2/source
[INFO] started tweaking crates.io crate svod-codegen 0.2.0
[INFO] finished tweaking crates.io crate svod-codegen 0.2.0
[INFO] tweaked toml for crates.io crate svod-codegen 0.2.0 written to /workspace/builds/worker-7-tc2/source/Cargo.toml
[INFO] validating manifest of crates.io crate svod-codegen 0.2.0 on toolchain 1.100.0-beta.1
[INFO] running `Command { std: CARGO_HOME="/workspace/cargo-home" RUSTUP_HOME="/workspace/rustup-home" "/workspace/cargo-home/bin/cargo" "+1.100.0-beta.1" "metadata" "--manifest-path" "Cargo.toml" "--no-deps", kill_on_drop: false }`
[INFO] crate crates.io crate svod-codegen 0.2.0 already has a lockfile, it will not be regenerated
[INFO] running `Command { std: CARGO_HOME="/workspace/cargo-home" RUSTUP_HOME="/workspace/rustup-home" "/workspace/cargo-home/bin/cargo" "+1.100.0-beta.1" "fetch" "--manifest-path" "Cargo.toml", kill_on_drop: false }`
[INFO] running `Command { std: "docker" "create" "-v" "/var/lib/crater-agent-workspace/builds/worker-7-tc2/source:/opt/rustwide/workdir:ro,Z" "-v" "/var/lib/crater-agent-workspace/builds/worker-7-tc2/target:/opt/rustwide/target:rw,Z" "-v" "/var/lib/crater-agent-workspace/cargo-home:/opt/rustwide/cargo-home:ro,Z" "-v" "/var/lib/crater-agent-workspace/rustup-home:/opt/rustwide/rustup-home:ro,Z" "-m" "1610612736" "--network" "none" "ghcr.io/rust-lang/crates-build-env/linux@sha256:77db811e55d90add9212f6832c23229e3de5f1c1f9905e8cb07319ea15df73ac" "sleep" "infinity", kill_on_drop: false }`
[INFO] [stdout] 2c77b06722d8a908026bfcb438b97adc842c6fc92fbd91244fc7d443de04deac
[INFO] running `Command { std: "docker" "start" "2c77b06722d8a908026bfcb438b97adc842c6fc92fbd91244fc7d443de04deac", kill_on_drop: false }`
[INFO] running `Command { std: "docker" "inspect" "2c77b06722d8a908026bfcb438b97adc842c6fc92fbd91244fc7d443de04deac", kill_on_drop: false }`
[INFO] running `Command { std: "docker" "exec" "-e" "SOURCE_DIR=/opt/rustwide/workdir" "-e" "CARGO_HOME=/opt/rustwide/cargo-home" "-e" "RUSTUP_HOME=/opt/rustwide/rustup-home" "-e" "CARGO_TARGET_DIR=/opt/rustwide/target" "-w" "/opt/rustwide/workdir" "--user" "0:0" "2c77b06722d8a908026bfcb438b97adc842c6fc92fbd91244fc7d443de04deac" "/opt/rustwide/cargo-home/bin/cargo" "+1.100.0-beta.1" "metadata" "--no-deps" "--format-version=1", kill_on_drop: false }`
[INFO] running `Command { std: "docker" "inspect" "2c77b06722d8a908026bfcb438b97adc842c6fc92fbd91244fc7d443de04deac", kill_on_drop: false }`
[INFO] running `Command { std: "docker" "exec" "-e" "SOURCE_DIR=/opt/rustwide/workdir" "-e" "CARGO_HOME=/opt/rustwide/cargo-home" "-e" "RUSTUP_HOME=/opt/rustwide/rustup-home" "-e" "CARGO_TARGET_DIR=/opt/rustwide/target" "-e" "CARGO_INCREMENTAL=0" "-e" "RUST_BACKTRACE=full" "-e" "RUSTFLAGS=--cap-lints=warn" "-e" "RUSTDOCFLAGS=--cap-lints=warn" "-w" "/opt/rustwide/workdir" "--user" "0:0" "2c77b06722d8a908026bfcb438b97adc842c6fc92fbd91244fc7d443de04deac" "/opt/rustwide/cargo-home/bin/cargo" "+1.100.0-beta.1" "build" "--frozen" "--message-format=json" "--release", kill_on_drop: false }`
[INFO] [stderr]    Compiling libc v0.2.190
[INFO] [stderr]    Compiling prettyplease v0.3.0
[INFO] [stderr]    Compiling clang-sys v1.9.1
[INFO] [stderr]    Compiling typenum v1.20.1
[INFO] [stderr]    Compiling regex-syntax v0.8.11
[INFO] [stderr]    Compiling memchr v2.8.3
[INFO] [stderr]    Compiling bindgen v0.73.2
[INFO] [stderr]    Compiling unicode-segmentation v1.13.3
[INFO] [stderr]    Compiling matrixmultiply v0.3.11
[INFO] [stderr]    Compiling syn v3.0.6
[INFO] [stderr]    Compiling syn v2.0.119
[INFO] [stderr]    Compiling parking_lot_core v0.8.6
[INFO] [stderr]    Compiling virtue v0.0.18
[INFO] [stderr]    Compiling cfg_aliases v0.2.2
[INFO] [stderr]    Compiling serde_json v1.0.151
[INFO] [stderr]    Compiling nix v0.31.3
[INFO] [stderr]    Compiling instant v0.1.13
[INFO] [stderr]    Compiling unty v0.0.4
[INFO] [stderr]    Compiling rawpointer v0.2.1
[INFO] [stderr]    Compiling const-oid v0.10.2
[INFO] [stderr]    Compiling crossbeam-epoch v0.9.21
[INFO] [stderr]    Compiling crc32fast v1.5.2
[INFO] [stderr]    Compiling num-integer v0.1.47
[INFO] [stderr]    Compiling num-complex v0.4.6
[INFO] [stderr]    Compiling convert_case v0.10.0
[INFO] [stderr]    Compiling option-ext v0.2.0
[INFO] [stderr]    Compiling xxhash-rust v0.8.19
[INFO] [stderr]    Compiling rustc-hash v2.1.3
[INFO] [stderr]    Compiling cpufeatures v0.3.1
[INFO] [stderr]    Compiling fxhash v0.2.1
[INFO] [stderr]    Compiling object v0.40.0
[INFO] [stderr]    Compiling libloading v0.9.0
[INFO] [stderr]    Compiling nom v7.1.3
[INFO] [stderr]    Compiling bincode_derive v2.0.1
[INFO] [stderr]    Compiling hybrid-array v0.4.15
[INFO] [stderr]    Compiling crypto-common v0.2.2
[INFO] [stderr]    Compiling block-buffer v0.12.1
[INFO] [stderr]    Compiling regex-automata v0.4.18
[INFO] [stderr]    Compiling ndarray v0.17.2
[INFO] [stderr]    Compiling cexpr v0.6.0
[INFO] [stderr]    Compiling seize v0.5.1
[INFO] [stderr]    Compiling dirs-sys v0.5.0
[INFO] [stderr]    Compiling memmap2 v0.9.11
[INFO] [stderr]    Compiling fs2 v0.4.3
[INFO] [stderr]    Compiling dirs v7.0.0
[INFO] [stderr]    Compiling digest v0.11.3
[INFO] [stderr]    Compiling papaya v0.2.5
[INFO] [stderr]    Compiling sha2 v0.11.0
[INFO] [stderr]    Compiling darling_core v0.24.1
[INFO] [stderr]    Compiling darling_core v0.21.3
[INFO] [stderr]    Compiling regex v1.13.1
[INFO] [stderr]    Compiling strum_macros v0.28.0
[INFO] [stderr]    Compiling derive_more-impl v2.1.1
[INFO] [stderr]    Compiling tracing-attributes v0.1.31
[INFO] [stderr]    Compiling snafu-derive v0.9.2
[INFO] [stderr]    Compiling serde_derive v1.0.229
[INFO] [stderr]    Compiling svod-macros v0.2.0
[INFO] [stderr]    Compiling tracing v0.1.44
[INFO] [stderr]    Compiling snafu v0.9.2
[INFO] [stderr]    Compiling darling_macro v0.24.1
[INFO] [stderr]    Compiling darling v0.24.1
[INFO] [stderr]    Compiling bon-macros v3.10.2
[INFO] [stderr]    Compiling derive_more v2.1.1
[INFO] [stderr]    Compiling strum v0.28.0
[INFO] [stderr]    Compiling darling_macro v0.21.3
[INFO] [stderr]    Compiling darling v0.21.3
[INFO] [stderr]    Compiling enumset_derive v0.15.0
[INFO] [stderr]    Compiling serde v1.0.229
[INFO] [stderr]    Compiling enumset v1.1.14
[INFO] [stderr]    Compiling smallvec v1.16.2
[INFO] [stderr]    Compiling bincode v2.0.1
[INFO] [stderr]    Compiling svod-dtype v0.2.0
[INFO] [stderr]    Compiling parking_lot_core v0.9.12
[INFO] [stderr]    Compiling parking_lot v0.12.5
[INFO] [stderr]    Compiling parking_lot v0.11.2
[INFO] [stderr]    Compiling sled v0.34.7
[INFO] [stderr]    Compiling svod-device v0.2.0
[INFO] [stderr]    Compiling bon v3.10.2
[INFO] [stderr]    Compiling svod-ir v0.2.0
[INFO] [stderr]    Compiling svod-schedule v0.2.0
[INFO] [stderr]    Compiling svod-codegen v0.2.0 (/opt/rustwide/workdir)
[INFO] [stderr]     Finished `release` profile [optimized] target(s) in 2m 23s
[INFO] running `Command { std: "docker" "inspect" "2c77b06722d8a908026bfcb438b97adc842c6fc92fbd91244fc7d443de04deac", kill_on_drop: false }`
[INFO] running `Command { std: "docker" "exec" "-e" "SOURCE_DIR=/opt/rustwide/workdir" "-e" "CARGO_HOME=/opt/rustwide/cargo-home" "-e" "RUSTUP_HOME=/opt/rustwide/rustup-home" "-e" "CARGO_TARGET_DIR=/opt/rustwide/target" "-e" "CARGO_INCREMENTAL=0" "-e" "RUST_BACKTRACE=full" "-e" "RUSTFLAGS=--cap-lints=warn" "-e" "RUSTDOCFLAGS=--cap-lints=warn" "-w" "/opt/rustwide/workdir" "--user" "0:0" "2c77b06722d8a908026bfcb438b97adc842c6fc92fbd91244fc7d443de04deac" "/opt/rustwide/cargo-home/bin/cargo" "+1.100.0-beta.1" "test" "--frozen" "--no-run" "--message-format=json" "--release", kill_on_drop: false }`
[INFO] [stderr]    Compiling test-case-core v3.4.0
[INFO] [stderr]    Compiling test-case-macros v3.4.0
[INFO] [stderr]    Compiling test-case v3.4.0
[INFO] [stderr]    Compiling svod-codegen v0.2.0 (/opt/rustwide/workdir)
[INFO] [stderr]     Finished `release` profile [optimized] target(s) in 17.39s
[INFO] running `Command { std: "docker" "inspect" "2c77b06722d8a908026bfcb438b97adc842c6fc92fbd91244fc7d443de04deac", kill_on_drop: false }`
[INFO] running `Command { std: "docker" "exec" "-e" "SOURCE_DIR=/opt/rustwide/workdir" "-e" "CARGO_HOME=/opt/rustwide/cargo-home" "-e" "RUSTUP_HOME=/opt/rustwide/rustup-home" "-e" "CARGO_TARGET_DIR=/opt/rustwide/target" "-e" "CARGO_INCREMENTAL=0" "-e" "RUST_BACKTRACE=full" "-e" "RUSTFLAGS=--cap-lints=warn" "-e" "RUSTDOCFLAGS=--cap-lints=warn" "-w" "/opt/rustwide/workdir" "--user" "0:0" "2c77b06722d8a908026bfcb438b97adc842c6fc92fbd91244fc7d443de04deac" "/opt/rustwide/cargo-home/bin/cargo" "+1.100.0-beta.1" "test" "--frozen" "--release", kill_on_drop: false }`
[INFO] [stderr]     Finished `release` profile [optimized] target(s) in 0.12s
[INFO] [stderr]      Running unittests src/lib.rs (/opt/rustwide/target/release/build/svod-codegen/be826287fb3a3d25/out/svod_codegen-be826287fb3a3d25)
[INFO] [stdout] 
[INFO] [stdout] running 345 tests
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_11_gfx950_fp8e4m3_float32_16_expects_none ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_15_dotted_k32_f16_is_gfx950_only ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_1_no_wmma_flavor_takes_bool ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_20_gfx950_float16_float32_64_expects_none ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_16_dotted_k32_bf16_is_gfx950_only ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_14_gfx950_bfloat16_float32_32_expects_some_llvm_amdgcn_mfma_f32_16x16x32_bf16_into_ ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_18_gfx942_has_no_scaled_k128_mfma ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_21_scaled_mfma_is_fp8_only ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_12_cdna4_keeps_the_unscaled_k32_fp8_mfma ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_17_gfx950_fp8e4m3_float32_128_expects_some_llvm_amdgcn_mfma_scale_f32_16x16x128_f8f6f4_into_ ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_13_gfx950_float16_float32_32_expects_some_llvm_amdgcn_mfma_f32_16x16x32_f16_into_ ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_19_gfx950_float16_float32_8_expects_none ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_25_gfx1100_float16_float32_16_expects_some_llvm_amdgcn_wmma_f32_16x16x16_f16_into_ ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_28_rdna4_has_no_native_fp8_wmma ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_29_gfx1151_fp8e4m3_float32_16_expects_none ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_26_gfx1201_float16_float32_16_expects_some_llvm_amdgcn_wmma_f32_16x16x16_f16_v8f32_v8f16_into_ ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_27_gfx1201_bfloat16_bfloat16_16_expects_some_llvm_amdgcn_wmma_bf16_16x16x16_bf16_v8i16_v8i16_into_ ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_22_gfx950_bfloat16_float32_8_expects_none ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_23_gfx950_bfloat16_float32_64_expects_none ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_24_gfx950_bfloat16_float32_128_expects_none ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_2_gfx942_bfloat16_float32_16_expects_some_llvm_amdgcn_mfma_f32_16x16x16bf16_1k_into_ ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_30_rdna_must_not_inherit_the_cdna_fp8_mfma ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_31_gfx1100_int8_int32_16_expects_some_llvm_amdgcn_wmma_i32_16x16x16_iu8_into_ ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_32_gfx1201_int8_int32_16_expects_some_llvm_amdgcn_wmma_i32_16x16x16_iu8_v8i32_v2i32_into_ ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_8_gfx942_fp8e4m3_float32_32_expects_some_llvm_amdgcn_mfma_f32_16x16x32_fp8_fp8_into_ ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_34_rdna_wmma_is_k16_only ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_36_the_accumulator_is_pinned_to_the_operand ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_3_gfx942_float16_float32_16_expects_some_llvm_amdgcn_mfma_f32_16x16x16f16_into_ ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_5_gfx942_float32_float32_8_expects_none ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_35_cdna_has_no_iu8_mfma_here ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_6_gfx942_float32_float32_16_expects_none ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_4_gfx942_float32_float32_4_expects_some_llvm_amdgcn_mfma_f32_16x16x4f32_into_ ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::rdna_int8_uses_packed_i32_wire_type::test_case_2_rdna4 ... ok
[INFO] [stdout] test llvm::common::gpu::tests::parse_special_axis_g_l_i ... ok
[INFO] [stdout] test llvm::common::types::tests::test_lcast ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::clang_flags_pin_the_ptx_isa_per_arch::test_case_4_hopper ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_7_gfx942_float32_float32_32_expects_none ... ok
[INFO] [stdout] test llvm::common::types::tests::test_lcast_index_type ... ok
[INFO] [stdout] test llvm::common::types::tests::test_ldt_ptr ... ok
[INFO] [stdout] test llvm::common::types::tests::test_ldt_scalar ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::clang_flags_pin_the_ptx_isa_per_arch::test_case_1_turing ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::clang_flags_pin_the_ptx_isa_per_arch::test_case_2_ampere ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::clang_flags_pin_the_ptx_isa_per_arch::test_case_3_ada_has_fp8 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::clang_flags_pin_the_ptx_isa_per_arch::test_case_5_blackwell_b100 ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_9_gfx942_fp8e5m2_float32_32_expects_some_llvm_amdgcn_mfma_f32_16x16x32_bf8_bf8_into_ ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::rdna_int8_uses_packed_i32_wire_type::test_case_1_rdna3 ... ok
[INFO] [stdout] test llvm::common::types::tests::test_lconst ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::clang_flags_pin_the_ptx_isa_per_arch::test_case_6_blackwell_b300 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::clang_flags_pin_the_ptx_isa_per_arch::test_case_7_blackwell_rtx_50 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::clang_flags_pin_the_ptx_isa_per_arch::test_case_8_blackwell_dgx_spark ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_33_the_int8_wmma_only_accumulates_in_i32 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_direct_global_axis_reads_the_block_index ... ok
[INFO] [stdout] test llvm::amd::wmma::tests::intrinsic_selection::test_case_10_gfx942_fp8e4m3_float32_16_expects_none ... ok
[INFO] [stdout] test llvm::common::types::tests::test_ldt_vector ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_casts_select::test_case_5_i32_to_f32 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_emits_kernel_abi_datalayout_and_launch_bound ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_rejects_f64_exp2_log2_without_decomposition::test_case_2_exp2 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_rejects_fp8_casts ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_rejects_undecomposed_transcendentals::test_case_1_sin ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_rejects_undecomposed_transcendentals::test_case_2_erf ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_log2_splits_vectors_per_lane ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_exp2_log2_lowering::test_case_3_f64_decomposes_polynomially ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_casts_select::test_case_6_f32_to_i64 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_casts_select::test_case_8_f32_to_bool_stores_a_byte ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_casts_select::test_case_3_f32_to_bf16_rounds_in_integers ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_exp2_log2_lowering::test_case_2_f16_widens_around_lg2 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_exp2_log2_lowering::test_case_1_f32 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_define_local_emits_addrspace3_module_global ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_casts_select::test_case_4_bf16_to_f32 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_register_buffers_use_generic_alloca ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_shared_memory_traffic_selects_shared_instructions ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_shfl_rejects_wide_values - should panic ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_shfl_rejects_shaped_values - should panic ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_casts_select::test_case_2_f16_to_f32 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_casts_select::test_case_1_f32_to_f16 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_rejects_f64_exp2_log2_without_decomposition::test_case_1_log2 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_barrier_emits_block_scope_fences_around_bar_sync ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvvm_builders_are_rejected_on_other_targets::test_case_2_amd ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvvm_builders_are_rejected_on_other_targets::test_case_1_cpu ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_casts_select::test_case_7_u8_to_f16 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_shfl_modes::test_case_1_idx ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_shfl_bfly_and_globaltimer_hoist_their_declares ... ok
[INFO] [stdout] test llvm::nvptx::smem::tests::cp_async_cg_is_16_bytes_only - should panic ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_shfl_moves_a_packed_half_pair_as_one_word ... ok
[INFO] [stdout] test llvm::nvptx::smem::tests::cp_async_rejects_a_shared_source - should panic ... ok
[INFO] [stdout] test llvm::nvptx::smem::tests::cp_async_rejects_a_global_destination - should panic ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_renders_for_older_capabilities ... ok
[INFO] [stdout] test llvm::nvptx::smem::tests::cp_async_ca_selects_the_sized_form::test_case_1_ca_4 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_special_emits_ctaid_and_tid_sregs ... ok
[INFO] [stdout] test llvm::nvptx::smem::tests::ldmatrix_rejects_narrow_fragments - should panic ... ok
[INFO] [stdout] test llvm::nvptx::smem::tests::ldmatrix_rejects_other_counts - should panic ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_shfl_modes::test_case_3_down ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_10_sm_8_6_bfloat16_float32_16_expects_some_llvm_nvvm_mma_m16n8k16_row_col_bf16_into_ ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_11_no_bf16_accumulator ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_12_sm_8_6_float32_float32_8_expects_some_llvm_nvvm_mma_m16n8k8_row_col_tf32_into_ ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_13_tf32_is_k8_only ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_14_sm_8_6_int8_int32_32_expects_some_llvm_nvvm_mma_m16n8k32_row_col_satfinite_s8_into_ ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_15_int8_is_k32_only ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_16_fp8_needs_ada ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_17_sm_8_9_fp8e4m3_float32_32_expects_some_llvm_nvvm_mma_m16n8k32_row_col_f32_e4m3_e4m3_f32_into_ ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_18_sm_8_9_fp8e5m2_float32_32_expects_some_llvm_nvvm_mma_m16n8k32_row_col_f32_e5m2_e5m2_f32_into_ ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_19_fp8_is_k32_only ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_1_turing_f16_k8 ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_20_open_ended_arch ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_21_f16_has_no_k32 ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_22_no_mma_form_takes_bool ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_2_turing_f16_acc_k8 ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_3_k16_needs_ampere ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_4_bf16_needs_ampere ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_5_tf32_needs_ampere ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_6_volta_has_no_m16n8_shape ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_7_sm_8_6_float16_float32_16_expects_some_llvm_nvvm_mma_m16n8k16_row_col_f32_f32_into_ ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_8_sm_8_6_float16_float16_16_expects_some_llvm_nvvm_mma_m16n8k16_row_col_f16_f16_into_ ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection::test_case_9_sm_8_6_float16_float32_8_expects_some_llvm_nvvm_mma_m16n8k8_row_col_f32_f32_into_ ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::mma_selection_requires_the_m16n8_shape ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::nvptx_wmma_below_the_shape_capability_is_rejected ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::nvptx_wmma_declaration_is_deduplicated_across_calls ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_shfl_widens_16bit_values::test_case_1_f16 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_shfl_moves_a_packed_byte_pair_as_one_word ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_shfl_widens_16bit_values::test_case_4_u16 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_shfl_modes::test_case_2_up_clamps_at_lane_0 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_shfl_modes::test_case_4_bfly ... ok
[INFO] [stdout] test llvm::nvptx::smem::tests::cp_async_16_casts_pointers_to_their_address_spaces ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing::test_case_3_cuda_81616_bf16_packs_i32_words ... ok
[INFO] [stdout] test llvm::nvptx::smem::tests::tile_prologue_compiles_to_ptx ... ok
[INFO] [stdout] test llvm::nvptx::smem::tests::cp_async_ca_selects_the_sized_form::test_case_3_ca_16 ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::nvptx_wmma_with_the_wrong_fragment_width_is_rejected ... ok
[INFO] [stdout] test llvm::text::tests::amd_barrier_emits_fence_and_s_barrier ... ok
[INFO] [stdout] test llvm::nvptx::smem::tests::ldmatrix_forms::test_case_2_x2_u32 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_shfl_widens_16bit_values::test_case_2_bf16 ... ok
[INFO] [stdout] test llvm::nvptx::smem::tests::ldmatrix_forms::test_case_3_x4_half_pairs ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing::test_case_8_cuda_81632_e5m2_on_ada ... FAILED
[INFO] [stdout] test llvm::nvptx::smem::tests::ldmatrix_forms::test_case_1_x1_i32 ... ok
[INFO] [stdout] test llvm::text::tests::amd_emits_kernel_abi_and_target_triple ... ok
[INFO] [stdout] test llvm::text::tests::amd_gfx1201_direct_fp8_wmma_is_rejected ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing::test_case_5_cuda_8168_f16_accumulator ... ok
[INFO] [stdout] test llvm::text::tests::amd_rejects_undecomposed_sin ... ok
[INFO] [stdout] test llvm::text::tests::amd_special_emits_workgroup_workitem_intrinsics ... ok
[INFO] [stdout] test llvm::text::tests::amd_uses_llvm_intrinsics_and_contract_only_float_flags::test_case_1_dtype_float16_llvm_exp2_f16_half_expects ... ok
[INFO] [stdout] test llvm::text::tests::amd_uses_llvm_intrinsics_and_contract_only_float_flags::test_case_3_dtype_float64_ocml_exp2_f64_double_expects ... ok
[INFO] [stdout] test llvm::text::tests::amd_uses_llvm_intrinsics_and_contract_only_float_flags::test_case_2_dtype_float32_llvm_exp2_f32_float_expects ... ok
[INFO] [stdout] test llvm::text::tests::amd_define_local_emits_addrspace3_module_global ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing::test_case_4_cuda_8168_f16_on_turing ... ok
[INFO] [stdout] test llvm::nvptx::smem::tests::ldmatrix_forms::test_case_4_x4_trans_bf16_pairs ... ok
[INFO] [stdout] test llvm::nvptx::smem::tests::cp_async_ca_selects_the_sized_form::test_case_2_ca_8 ... ok
[INFO] [stdout] test llvm::nvptx::ops::tests::nvptx_shfl_widens_16bit_values::test_case_3_i16 ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing::test_case_9_int8_satfinite ... ok
[INFO] [stdout] test llvm::text::tests::amd_cdna_ordinary_fp8_alu_widens_without_changing_storage::test_case_1_amdarch_gfx942_dtype_fp8e4m3_expects ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing::test_case_7_cuda_81632_e4m3_on_ada ... FAILED
[INFO] [stdout] test llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing::test_case_1_cuda_81616_f16_to_f32 ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing::test_case_2_cuda_81616_f16_accumulator ... ok
[INFO] [stdout] test llvm::text::tests::custom_renders_an_ssa_rhs_and_customi_inlines_into_its_consumer ... ok
[INFO] [stdout] test llvm::text::tests::custom_void_hoists_and_deduplicates_declares_into_the_module_prefix ... ok
[INFO] [stdout] test llvm::text::tests::llvm_float_comparisons_use_ordered_predicates_matching_c ... ok
[INFO] [stdout] test llvm::text::tests::llvm_rejects_malformed_range_nesting::test_case_1_range_never_closed ... ok
[INFO] [stdout] test llvm::text::tests::llvm_rejects_malformed_range_nesting::test_case_2_end_targets_an_outer_range ... ok
[INFO] [stdout] test llvm::text::tests::amd_cdna_ordinary_fp8_alu_widens_without_changing_storage::test_case_3_amdarch_gfx950_dtype_fp8e4m3_expects ... ok
[INFO] [stdout] test llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing::test_case_6_cuda_8168_tf32_passes_raw_f32_words ... ok
[INFO] [stdout] test llvm::text::tests::amd_wmma_operand_packing::test_case_2_rdna3_bf16_goes_on_the_wire_as_i16 ... ok
[INFO] [stdout] test llvm::text::tests::amd_wmma_operand_packing::test_case_10_cdna4_bf16_k32_takes_native_bfloat_operands ... ok
[INFO] [stdout] test llvm::text::tests::amd_cdna_ordinary_fp8_alu_widens_without_changing_storage::test_case_4_amdarch_gfx950_dtype_fp8e5m2_expects ... ok
[INFO] [stdout] test llvm::text::tests::llvm_shaped_values_assemble::test_case_1_stack_of_loaded_index_lanes ... ok
[INFO] [stdout] test llvm::text::tests::llvm_shaped_values_assemble::test_case_2_cast_of_shaped_stack ... ok
[INFO] [stdout] test llvm::text::tests::llvm_signature_uses_canonical_mixed_param_slot_order ... ok
[INFO] [stdout] test llvm::text::tests::test_simple_add ... ok
[INFO] [stdout] test test::unit::c::c_constants_consume_committed_values_and_fp8_bits ... ok
[INFO] [stdout] test test::unit::c::c_qualifies_only_volatile_buffer_parameters ... ok
[INFO] [stdout] test llvm::text::tests::amd_wmma_operand_packing::test_case_5_gfx1200_uses_llvm_overloaded_vector_suffixes ... ok
[INFO] [stdout] test llvm::text::tests::amd_wmma_operand_packing::test_case_1_rdna3_f16_to_f32 ... ok
[INFO] [stdout] test llvm::text::tests::llvm_shaped_values_assemble::test_case_4_gated_load_with_shaped_alt ... ok
[INFO] [stdout] test test::unit::c::c_renders_range_and_reduce_structure::test_case_2_sum_accumulates_from_the_identity ... ok
[INFO] [stdout] test test::unit::c::c_signature_uses_canonical_mixed_param_slot_order ... ok
[INFO] [stdout] test llvm::text::tests::amd_wmma_operand_packing::test_case_3_rdna3_bf16_accumulator_bitcasts_the_result_back ... ok
[INFO] [stdout] test test::unit::c::clang_materializes_tinygrad_ssa_boundaries ... ok
[INFO] [stdout] test llvm::text::tests::amd_wmma_operand_packing::test_case_6_gfx1201_uses_llvm_overloaded_vector_suffixes ... ok
[INFO] [stdout] test llvm::text::tests::amd_wmma_operand_packing::test_case_11_cdna4_scaled_fp8_uses_i32_vectors_and_scale_immediates ... ok
[INFO] [stdout] test llvm::text::tests::llvm_shaped_values_assemble::test_case_7_store_vector_through_scalar_index ... ok
[INFO] [stdout] test llvm::text::tests::amd_wmma_operand_packing::test_case_4_rdna3_int8_packs_into_i32_words_with_signedness_flags ... ok
[INFO] [stdout] test llvm::text::tests::amd_cdna_ordinary_fp8_alu_widens_without_changing_storage::test_case_2_amdarch_gfx942_dtype_fp8e5m2_expects ... ok
[INFO] [stdout] test test::unit::c::clang_preserves_shape_width_across_materialized_values_and_store_aliases ... ok
[INFO] [stdout] test test::unit::c::clang_preserves_address_casts_as_pointers ... ok
[INFO] [stdout] test test::unit::c::c_renders_range_and_reduce_structure::test_case_1_loop_over_a_concrete_range ... ok
[INFO] [stdout] test test::unit::c::clang_narrows_promoted_arithmetic::test_case_1_int8_add ... ok
[INFO] [stdout] test test::unit::c::custom_template_errors::test_case_1_index_past_the_last_operand ... ok
[INFO] [stdout] test test::unit::c::custom_template_errors::test_case_2_unclosed_brace ... ok
[INFO] [stdout] test test::unit::c::clang_vector_alignment_rounds_down_like_tinygrad ... ok
[INFO] [stdout] test test::unit::c::custom_template_errors::test_case_3_automatic_and_manual_indices_mixed ... ok
[INFO] [stdout] test test::unit::c::c_renders_range_and_reduce_structure::test_case_4_min_starts_from_inf_and_uses_fminf ... ok
[INFO] [stdout] test test::unit::c::grouped_shrink_renders_single_vector_load_and_store::test_case_2_f16 ... ok
[INFO] [stdout] test test::unit::c::test_custom_statement_is_materialized ... ok
[INFO] [stdout] test test::unit::c::test_customi_is_inline_and_formats_placeholders ... ok
[INFO] [stdout] test test::unit::c::test_gated_load_emits_conditional_dereference ... ok
[INFO] [stdout] test test::unit::c::test_getaddr_must_be_resolved_before_codegen ... ok
[INFO] [stdout] test test::unit::c::clang_stack_packs_loaded_index_lanes ... ok
[INFO] [stdout] test test::unit::c::test_multi_index_requires_linearization ... ok
[INFO] [stdout] test test::unit::c::test_render_linear_input_succeeds ... ok
[INFO] [stdout] test test::unit::c::test_render_rejects_fnuz_without_fallback ... ok
[INFO] [stdout] test test::unit::c::test_render_rejects_non_linear_inputs ... ok
[INFO] [stdout] test test::unit::llvm::grouped_shrink_renders_single_vector_load_and_store ... ok
[INFO] [stdout] test test::unit::llvm::llvm_constants_use_committed_storage_bits ... ok
[INFO] [stdout] test test::unit::llvm::custom_renders_a_typed_rhs_and_customi_inlines_into_its_consumer ... ok
[INFO] [stdout] test test::unit::llvm::llvm_gated_load_phis_the_shaped_load_type::test_case_2_grouped_gated_load ... ok
[INFO] [stdout] test test::unit::llvm::shaped_reg_index_renders_as_memory_load::test_case_1_direct ... ok
[INFO] [stdout] test test::unit::llvm::shaped_reg_index_renders_as_memory_load::test_case_2_through_after ... ok
[INFO] [stdout] test test::unit::llvm::test_multi_index_requires_linearization ... ok
[INFO] [stdout] test llvm::text::tests::amd_wmma_operand_packing::test_case_7_gfx1201_int8_packs_eight_lanes_into_two_i32_words ... ok
[INFO] [stdout] test test::unit::llvm::test_range_end_basic ... ok
[INFO] [stdout] test test::unit::llvm::test_render_rejects_non_linear_inputs ... ok
[INFO] [stdout] test test::unit::llvm::test_render_linear_input_succeeds ... ok
[INFO] [stdout] test llvm::text::tests::llvm_shaped_values_assemble::test_case_5_gated_load_with_scalar_alt ... ok
[INFO] [stdout] test test::unit::metal::metal_address_casts_carry_the_address_space ... ok
[INFO] [stdout] test test::unit::metal::metal_barrier_emits_threadgroup_barrier ... ok
[INFO] [stdout] test test::unit::llvm::volatile_scalar_and_grouped_memory_accesses_render_explicitly ... ok
[INFO] [stdout] test llvm::text::tests::amd_wmma_operand_packing::test_case_9_cdna4_fp8_packs_eight_lanes_into_one_i64 ... ok
[INFO] [stdout] test test::unit::c::clang_narrows_promoted_arithmetic::test_case_2_uint16_add ... ok
[INFO] [stdout] test test::unit::metal::metal_constants::test_case_12_constvalue_uint_5_dtype_uint32_expects_5u_ ... ok
[INFO] [stdout] test test::unit::metal::metal_constants::test_case_13_constvalue_int_5_dtype_int32_expects_5_ ... ok
[INFO] [stdout] test test::unit::c::clang_hoists_addresses_that_escape_their_loop ... ok
[INFO] [stdout] test test::unit::metal::metal_constants::test_case_1_constvalue_float_3_2_dtype_float32_expects_3_2e0f_ ... ok
[INFO] [stdout] test test::unit::metal::metal_constants::test_case_2_constvalue_float_1_0_dtype_float16_expects_half_1_0f_ ... ok
[INFO] [stdout] test test::unit::metal::metal_constants::test_case_11_constvalue_uint_5_dtype_uint64_expects_5ul_ ... ok
[INFO] [stdout] test test::unit::metal::metal_constants::test_case_10_constvalue_int_5_dtype_index_expects_5l_ ... ok
[INFO] [stdout] test test::unit::metal::metal_constants::test_case_5_constvalue_float_f64_neg_infinity_dtype_float32_expects_infinity_ ... ok
[INFO] [stdout] test test::unit::c::c_renders_range_and_reduce_structure::test_case_3_max_starts_from_inf_and_uses_fmaxf ... ok
[INFO] [stdout] test test::unit::metal::metal_constants::test_case_8_constvalue_float_f64_nan_dtype_bfloat16_expects_bfloat_nan_ ... ok
[INFO] [stdout] test test::unit::metal::metal_constants::test_case_3_constvalue_float_1_0_dtype_bfloat16_expects_bfloat_1_0f_ ... ok
[INFO] [stdout] test test::unit::metal::metal_constants::test_case_6_constvalue_float_f64_infinity_dtype_float16_expects_half_infinity_ ... ok
[INFO] [stdout] test test::unit::metal::metal_constants::test_case_7_constvalue_float_f64_nan_dtype_float32_expects_nan_ ... ok
[INFO] [stdout] test test::unit::metal::metal_golden_threadgroup_reduction ... ok
[INFO] [stdout] test test::unit::metal::metal_local_buffer_is_threadgroup_memory ... ok
[INFO] [stdout] test test::unit::metal::metal_golden_elementwise_add ... ok
[INFO] [stdout] test test::unit::c::reduce_without_ranges_renders_as_a_plain_value ... ok
[INFO] [stdout] test test::unit::llvm::llvm_gated_load_phis_the_shaped_load_type::test_case_1_scalar_gated_load ... ok
[INFO] [stdout] test test::unit::metal::clang_dialect_output_is_unchanged ... ok
[INFO] [stdout] test test::unit::metal::metal_constants::test_case_4_constvalue_float_f64_infinity_dtype_float32_expects_infinity_ ... ok
[INFO] [stdout] test test::unit::metal::metal_math_functions::test_case_1_sqrt_f32 ... ok
[INFO] [stdout] test test::unit::metal::metal_math_functions::test_case_6_bfloat_promotes ... ok
[INFO] [stdout] test test::unit::metal::metal_bitcast_uses_as_type ... ok
[INFO] [stdout] test test::unit::c::clang_store_width_follows_the_stored_value ... ok
[INFO] [stdout] test test::unit::metal::metal_math_functions::test_case_4_log2_f32 ... ok
[INFO] [stdout] test test::unit::metal::metal_prelude_has_no_clang_artifacts ... ok
[INFO] [stdout] test test::unit::metal::metal_narrow_int_arithmetic_casts_back ... ok
[INFO] [stdout] test test::unit::metal::metal_math_functions::test_case_2_sin_uses_precise ... ok
[INFO] [stdout] test test::unit::metal::metal_reduce_structure::test_case_2_max ... ok
[INFO] [stdout] test test::unit::metal::metal_math_functions::test_case_3_exp2_f32 ... ok
[INFO] [stdout] test test::unit::metal::metal_rejects_unsupported_dtypes::test_case_1_float64 ... ok
[INFO] [stdout] test test::unit::metal::metal_math_functions::test_case_8_arithmetic ... ok
[INFO] [stdout] test test::unit::metal::metal_rejects_unsupported_dtypes::test_case_2_fp8 ... ok
[INFO] [stdout] test test::unit::metal::metal_rejects_unsupported_dtypes::test_case_3_fnuz ... ok
[INFO] [stdout] test test::unit::metal::metal_rejects_unsupported_dtypes::test_case_4_wide_vector ... ok
[INFO] [stdout] test test::unit::metal::metal_math_functions::test_case_5_half_rounds_back ... ok
[INFO] [stdout] test test::unit::metal::metal_reg_buffer_is_a_thread_array ... ok
[INFO] [stdout] test test::unit::metal::metal_scalar_type_names::test_case_10_dtype_float16_half_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_scalar_type_names::test_case_12_dtype_float32_float_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_scalar_type_names::test_case_1_dtype_bool_bool_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_scalar_type_names::test_case_11_dtype_bfloat16_bfloat_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_math_functions::test_case_7_rsqrt ... ok
[INFO] [stdout] test test::unit::metal::metal_scalar_type_names::test_case_3_dtype_uint8_uchar_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_scalar_type_names::test_case_5_dtype_uint16_ushort_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_scalar_type_names::test_case_2_dtype_int8_char_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_reduce_structure::test_case_3_min ... ok
[INFO] [stdout] test test::unit::metal::metal_scalar_type_names::test_case_4_dtype_int16_short_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_scalar_type_names::test_case_7_dtype_uint32_uint_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_simd_shuffle_rejects_64_bit::test_case_1_i64 - should panic ... ok
[INFO] [stdout] test test::unit::metal::metal_simd_shuffle_rejects_64_bit::test_case_2_u64 - should panic ... ok
[INFO] [stdout] test test::unit::c::clang_narrows_promoted_arithmetic::test_case_3_half_sqrt ... ok
[INFO] [stdout] test test::unit::metal::metal_simd_shuffle_renders::test_case_3_f16_shuffle ... ok
[INFO] [stdout] test llvm::text::tests::amd_wmma_operand_packing::test_case_8_cdna3_fp8_packs_eight_lanes_into_one_i64 ... ok
[INFO] [stdout] test test::unit::metal::metal_signature_uses_device_and_constant_qualifiers ... ok
[INFO] [stdout] test test::unit::metal::metal_scalar_type_names::test_case_8_dtype_int64_long_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_scalar_type_names::test_case_6_dtype_int32_int_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_hoists_addresses_with_their_address_space ... ok
[INFO] [stdout] test test::unit::metal::metal_simd_shuffle_xor_renders ... ok
[INFO] [stdout] test test::unit::metal::metal_simd_shuffle_renders::test_case_1_f32_shuffle ... ok
[INFO] [stdout] test test::unit::metal::metal_simd_shuffle_xor_rejects_64_bit - should panic ... ok
[INFO] [stdout] test test::unit::metal::metal_simd_shuffle_renders::test_case_2_i32_shuffle ... ok
[INFO] [stdout] test test::unit::metal::metal_special_axes::test_case_1_gidx0_int_gidx0_gid_x_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_scalar_type_names::test_case_9_dtype_uint64_ulong_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_stack_uses_vector_constructor ... ok
[INFO] [stdout] test test::unit::metal::metal_special_with_malformed_name_is_invalid_graph ... ok
[INFO] [stdout] test test::unit::metal::metal_special_axes::test_case_5_lidx1_int_lidx1_lid_y_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_reduce_structure::test_case_1_sum ... ok
[INFO] [stdout] test test::unit::metal::metal_special_axes::test_case_2_gidx1_int_gidx1_gid_y_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_special_axes::test_case_4_lidx0_int_lidx0_lid_x_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_volatile_buffer_keeps_the_qualifier ... ok
[INFO] [stdout] test test::unit::metal::metal_vector_types_are_native::test_case_2_dtype_float16_4_half4_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_special_axes::test_case_7_idx0_int_idx0_gid_x_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_wmma_helpers_are_deduped_per_shape ... ok
[INFO] [stdout] test test::unit::metal::metal_wmma_lowers_to_simdgroup_matrix::test_case_2_dtype_float16_dtype_float32_half_float_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_wmma_lowers_to_simdgroup_matrix::test_case_4_dtype_bfloat16_dtype_float32_bfloat_float_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_wmma_lowers_to_simdgroup_matrix::test_case_5_dtype_bfloat16_dtype_bfloat16_bfloat_bfloat_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_wmma_rejects_non_simdgroup_shapes ... ok
[INFO] [stdout] test test::unit::c::clang_narrows_promoted_arithmetic::test_case_4_bfloat16_sqrt ... ok
[INFO] [stdout] test test::unit::metal::metal_special_axes::test_case_6_lidx2_int_lidx2_lid_z_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_vector_types_are_native::test_case_4_dtype_uint8_4_uchar4_expects ... ok
[INFO] [stdout] test test::unit::program_pipeline::abi_slots_are_assigned_after_authored_storage::test_case_1_one_global_reserves_slot_zero ... ok
[INFO] [stdout] test test::unit::program_pipeline::abi_slots_are_assigned_after_authored_storage::test_case_4_local_and_register_storage_stay_out_of_the_scalar_namespace ... ok
[INFO] [stdout] test test::unit::program_pipeline::abi_slots_are_assigned_after_authored_storage::test_case_2_scalars_follow_dense_globals_in_walk_order ... ok
[INFO] [stdout] test test::unit::program_pipeline::abi_slots_are_assigned_after_authored_storage::test_case_3_sparse_storage_slots_keep_their_gaps ... ok
[INFO] [stdout] test test::unit::metal::metal_vector_types_are_native::test_case_1_dtype_float32_4_float4_expects ... ok
[INFO] [stdout] test test::unit::program_pipeline::binary_stages_from_another_signature_or_compiler_are_rejected ... ok
[INFO] [stdout] test test::unit::metal::metal_wmma_lowers_to_simdgroup_matrix::test_case_1_dtype_float32_dtype_float32_float_float_expects ... ok
[INFO] [stdout] test test::unit::program_pipeline::do_render_rejects_renderer_abi_disagreeing_with_program_info::test_case_1_storage_count_with_no_descriptors ... ok
[INFO] [stdout] test test::unit::metal::metal_special_axes::test_case_3_gidx2_int_gidx2_gid_z_expects ... ok
[INFO] [stdout] test test::unit::program_pipeline::final_sink_verification_cannot_be_bypassed_by_a_staged_linear ... ok
[INFO] [stdout] test test::unit::program_pipeline::do_render_rejects_renderer_abi_disagreeing_with_program_info::test_case_2_same_count_descriptors_in_the_wrong_order ... ok
[INFO] [stdout] test test::unit::program_pipeline::conflicting_and_unassigned_abi_slots_are_typed_errors ... ok
[INFO] [stdout] test test::unit::metal::metal_vector_types_are_native::test_case_3_dtype_int32_2_int2_expects ... ok
[INFO] [stdout] test test::unit::program_pipeline::do_compile_rejects_malformed_stages::test_case_1_nothing_to_compile ... ok
[INFO] [stdout] test test::unit::program_pipeline::do_compile_rejects_malformed_stages::test_case_2_binary_stage_is_not_a_programbinary ... ok
[INFO] [stdout] test test::unit::program_pipeline::do_compile_rejects_malformed_stages::test_case_3_source_stage_is_empty ... ok
[INFO] [stdout] test test::unit::metal::metal_vector_constant_uses_constructor_and_store_width_follows_value ... ok
[INFO] [stdout] test test::unit::program_pipeline::gated_store_linearizes_to_a_verified_if_block ... ok
[INFO] [stdout] test test::unit::program_pipeline::number_params_uses_walk_after_control_flow_insertion ... ok
[INFO] [stdout] test test::unit::program_pipeline::get_program_rejects_non_advanceable_input::test_case_1_source_staged_without_linear ... ok
[INFO] [stdout] test test::unit::program_pipeline::prebuilt_program_rejects_descriptor_equivalent_var_forgery::test_case_1_bounds ... ok
[INFO] [stdout] test test::unit::program_pipeline::get_program_rejects_non_advanceable_input::test_case_2_bare_sink ... ok
[INFO] [stdout] test test::unit::program_pipeline::opaque_function_formals_stay_out_of_the_outer_abi ... ok
[INFO] [stdout] test test::unit::program_pipeline::prebuilt_program_rejects_descriptor_equivalent_var_forgery::test_case_3_axis ... ok
[INFO] [stdout] test test::unit::program_pipeline::prebuilt_program_rejects_descriptor_equivalent_var_forgery::test_case_2_multiple_of ... ok
[INFO] [stdout] test test::unit::program_pipeline::isa_selection_is_bottom_up_after_program_info_and_renders_program_source ... ok
[INFO] [stdout] test test::unit::program_pipeline::get_program_advances_any_staged_program_to_a_reusable_binary ... ok
[INFO] [stdout] test test::unit::program_pipeline::prebuilt_program_target_must_match_renderer ... ok
[INFO] [stdout] test test::unit::program_pipeline::prebuilt_program_accepts_semantically_identical_nonidentical_var ... ok
[INFO] [stdout] test test::unit::program_pipeline::symbolic_program_render_compile_and_runtime_binding_share_canonical_abi ... ok
[INFO] [stdout] test test::unit::program_pipeline::renderers_emit_the_sparse_abi_verbatim ... ok
[INFO] [stdout] test test::unit::program_pipeline::test_do_render_uses_linear_stage_input ... ok
[INFO] [stdout] test test::unit::program_pipeline::test_hand_lowered_final_rewrite_stays_invalid_free_through_linearize ... ok
[INFO] [stdout] test test::unit::program_pipeline::reused_param_is_one_abi_argument ... ok
[INFO] [stdout] test test::unit::program_pipeline::test_do_compile_reuses_existing_binary_stage ... ok
[INFO] [stdout] test test::unit::program_pipeline::test_do_render_rejects_program_with_existing_binary_stage ... ok
[INFO] [stdout] test test::unit::program_pipeline::semantic_stage_identity_defeats_preinterned_children_and_parent_programs ... ok
[INFO] [stdout] test test::unit::program_pipeline::unnamed_scalar_and_non_param_program_info_var_are_typed_errors ... ok
[INFO] [stdout] test test::unit::program_pipeline::source_stages_without_a_matching_identity_are_rejected ... ok
[INFO] [stdout] test test::unit::program_pipeline::test_structured_symbolic_name_is_sanitized_at_renderer_boundary ... ok
[INFO] [stdout] test test::unit::program_pipeline::repeated_program_construction_has_identical_slots_and_identity ... ok
[INFO] [stdout] test test::unit::program_pipeline::test_program_pipeline_sets_all_stages ... ok
[INFO] [stdout] test llvm::text::tests::llvm_shaped_values_assemble::test_case_6_index_without_indices ... ok
[INFO] [stdout] test llvm::text::tests::llvm_shaped_values_assemble::test_case_3_index_into_shrink ... ok
[INFO] [stdout] test test::unit::program_pipeline::test_structured_custom_name_wins_over_optimizer_shape_name ... ok
[INFO] [stdout] test test::unit::c::grouped_shrink_renders_single_vector_load_and_store::test_case_1_f32 ... ok
[INFO] [stdout] test test::unit::metal::metal_constants::test_case_9_constvalue_int_5_dtype_int64_expects_5l_ ... ok
[INFO] [stdout] test test::unit::metal::metal_special_axes::test_case_8_idx1_int_idx1_gid_y_expects ... ok
[INFO] [stdout] test test::unit::metal::metal_local_vector_access_carries_the_address_space ... ok
[INFO] [stdout] test test::unit::metal::metal_wmma_lowers_to_simdgroup_matrix::test_case_3_dtype_float16_dtype_float16_half_half_expects ... ok
[INFO] [stderr] error: test failed, to rerun pass `--lib`
[INFO] [stdout] 
[INFO] [stdout] failures:
[INFO] [stdout] 
[INFO] [stdout] ---- llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing::test_case_8_cuda_81632_e5m2_on_ada stdout ----
[INFO] [stdout] 
[INFO] [stdout] thread 'llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing::test_case_8_cuda_81632_e5m2_on_ada' (2543) panicked at src/llvm/nvptx/../../test/unit/llvm_nvptx_ops.rs:76:5:
[INFO] [stdout] PTX carries an unresolved intrinsic call:
[INFO] [stdout] //
[INFO] [stdout] // Generated by LLVM NVPTX Back-End
[INFO] [stdout] //
[INFO] [stdout] 
[INFO] [stdout] .version 8.4
[INFO] [stdout] .target sm_89
[INFO] [stdout] .address_size 64
[INFO] [stdout] 
[INFO] [stdout] 	// .globl	nvptx_wmma              // -- Begin function nvptx_wmma
[INFO] [stdout] .extern .func  (.param .align 4 .b8 func_retval0[16]) llvm.nvvm.mma.m16n8k32.row.col.f32.e5m2.e5m2.f32
[INFO] [stdout] (
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e5m2.e5m2.f32_param_0,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e5m2.e5m2.f32_param_1,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e5m2.e5m2.f32_param_2,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e5m2.e5m2.f32_param_3,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e5m2.e5m2.f32_param_4,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e5m2.e5m2.f32_param_5,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e5m2.e5m2.f32_param_6,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e5m2.e5m2.f32_param_7,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e5m2.e5m2.f32_param_8,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e5m2.e5m2.f32_param_9
[INFO] [stdout] )
[INFO] [stdout] ;
[INFO] [stdout]                                         // @nvptx_wmma
[INFO] [stdout] .visible .entry nvptx_wmma(
[INFO] [stdout] 	.param .u64 .ptr .align 32 nvptx_wmma_param_0,
[INFO] [stdout] 	.param .u64 .ptr .align 32 nvptx_wmma_param_1,
[INFO] [stdout] 	.param .u64 .ptr .align 32 nvptx_wmma_param_2,
[INFO] [stdout] 	.param .u64 .ptr .align 32 nvptx_wmma_param_3
[INFO] [stdout] )
[INFO] [stdout] .maxntid 1
[INFO] [stdout] {
[INFO] [stdout] 	.reg .b32 	%r<13>;
[INFO] [stdout] 	.reg .b64 	%rd<9>;
[INFO] [stdout] 
[INFO] [stdout] // %bb.0:                               // %entry
[INFO] [stdout] 	ld.param.b64 	%rd1, [nvptx_wmma_param_0];
[INFO] [stdout] 	cvta.to.global.u64 	%rd2, %rd1;
[INFO] [stdout] 	ld.param.b64 	%rd3, [nvptx_wmma_param_1];
[INFO] [stdout] 	cvta.to.global.u64 	%rd4, %rd3;
[INFO] [stdout] 	ld.param.b64 	%rd5, [nvptx_wmma_param_2];
[INFO] [stdout] 	cvta.to.global.u64 	%rd6, %rd5;
[INFO] [stdout] 	ld.param.b64 	%rd7, [nvptx_wmma_param_3];
[INFO] [stdout] 	cvta.to.global.u64 	%rd8, %rd7;
[INFO] [stdout] 	ld.global.nc.b32 	%r1, [%rd6];
[INFO] [stdout] 	ld.global.nc.b8 	%r2, [%rd2];
[INFO] [stdout] 	prmt.b32 	%r3, %r2, %r2, 0x3340U;
[INFO] [stdout] 	prmt.b32 	%r4, %r3, %r3, 0x5410U;
[INFO] [stdout] 	ld.global.nc.b8 	%r5, [%rd4];
[INFO] [stdout] 	prmt.b32 	%r6, %r5, %r5, 0x3340U;
[INFO] [stdout] 	prmt.b32 	%r7, %r6, %r6, 0x5410U;
[INFO] [stdout] 	{ // callseq 0, 0
[INFO] [stdout] 	.param .b32 	param0;
[INFO] [stdout] 	st.param.b32 	[param0], %r4;
[INFO] [stdout] 	.param .b32 	param1;
[INFO] [stdout] 	st.param.b32 	[param1], %r4;
[INFO] [stdout] 	.param .b32 	param2;
[INFO] [stdout] 	st.param.b32 	[param2], %r4;
[INFO] [stdout] 	.param .b32 	param3;
[INFO] [stdout] 	st.param.b32 	[param3], %r4;
[INFO] [stdout] 	.param .b32 	param4;
[INFO] [stdout] 	st.param.b32 	[param4], %r7;
[INFO] [stdout] 	.param .b32 	param5;
[INFO] [stdout] 	st.param.b32 	[param5], %r7;
[INFO] [stdout] 	.param .b32 	param6;
[INFO] [stdout] 	st.param.b32 	[param6], %r1;
[INFO] [stdout] 	.param .b32 	param7;
[INFO] [stdout] 	st.param.b32 	[param7], %r1;
[INFO] [stdout] 	.param .b32 	param8;
[INFO] [stdout] 	st.param.b32 	[param8], %r1;
[INFO] [stdout] 	.param .b32 	param9;
[INFO] [stdout] 	st.param.b32 	[param9], %r1;
[INFO] [stdout] 	.param .align 4 .b8 	retval0[16];
[INFO] [stdout] 	call.uni (retval0), llvm.nvvm.mma.m16n8k32.row.col.f32.e5m2.e5m2.f32, (param0, param1, param2, param3, param4, param5, param6, param7, param8, param9);
[INFO] [stdout] 	ld.param.b32 	%r8, [retval0];
[INFO] [stdout] 	ld.param.b32 	%r9, [retval0+4];
[INFO] [stdout] 	ld.param.b32 	%r10, [retval0+8];
[INFO] [stdout] 	ld.param.b32 	%r11, [retval0+12];
[INFO] [stdout] 	} // callseq 0
[INFO] [stdout] 	st.global.b32 	[%rd8], %r11;
[INFO] [stdout] 	ret;
[INFO] [stdout]                                         // -- End function
[INFO] [stdout] }
[INFO] [stdout] 
[INFO] [stdout] --- IR ---
[INFO] [stdout] ; ModuleID = 'nvptx_wmma'
[INFO] [stdout] source_filename = "nvptx_wmma"
[INFO] [stdout] target datalayout = "e-p6:32:32-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64"
[INFO] [stdout] target triple = "nvptx64-nvidia-cuda"
[INFO] [stdout] 
[INFO] [stdout] declare { float, float, float, float } @llvm.nvvm.mma.m16n8k32.row.col.f32.e5m2.e5m2.f32(i32, i32, i32, i32, i32, i32, float, float, float, float)
[INFO] [stdout] 
[INFO] [stdout] define ptx_kernel void @nvptx_wmma(ptr noalias align 32 %data0, ptr noalias align 32 %data1, ptr noalias align 32 %data2, ptr noalias align 32 %data3) #0 {
[INFO] [stdout] entry:
[INFO] [stdout] 
[INFO] [stdout]   %v0 = getelementptr inbounds i8, ptr %data0, i32 0
[INFO] [stdout]   %v1 = load i8, ptr %v0
[INFO] [stdout]   %v2 = getelementptr inbounds i8, ptr %data1, i32 0
[INFO] [stdout]   %v3 = load i8, ptr %v2
[INFO] [stdout]   %v4 = getelementptr inbounds float, ptr %data2, i32 0
[INFO] [stdout]   %v5 = load float, ptr %v4
[INFO] [stdout]   %v6.v0 = insertelement <16 x i8> poison, i8 %v1, i32 0
[INFO] [stdout]   %v6.v1 = insertelement <16 x i8> %v6.v0, i8 %v1, i32 1
[INFO] [stdout]   %v6.v2 = insertelement <16 x i8> %v6.v1, i8 %v1, i32 2
[INFO] [stdout]   %v6.v3 = insertelement <16 x i8> %v6.v2, i8 %v1, i32 3
[INFO] [stdout]   %v6.v4 = insertelement <16 x i8> %v6.v3, i8 %v1, i32 4
[INFO] [stdout]   %v6.v5 = insertelement <16 x i8> %v6.v4, i8 %v1, i32 5
[INFO] [stdout]   %v6.v6 = insertelement <16 x i8> %v6.v5, i8 %v1, i32 6
[INFO] [stdout]   %v6.v7 = insertelement <16 x i8> %v6.v6, i8 %v1, i32 7
[INFO] [stdout]   %v6.v8 = insertelement <16 x i8> %v6.v7, i8 %v1, i32 8
[INFO] [stdout]   %v6.v9 = insertelement <16 x i8> %v6.v8, i8 %v1, i32 9
[INFO] [stdout]   %v6.v10 = insertelement <16 x i8> %v6.v9, i8 %v1, i32 10
[INFO] [stdout]   %v6.v11 = insertelement <16 x i8> %v6.v10, i8 %v1, i32 11
[INFO] [stdout]   %v6.v12 = insertelement <16 x i8> %v6.v11, i8 %v1, i32 12
[INFO] [stdout]   %v6.v13 = insertelement <16 x i8> %v6.v12, i8 %v1, i32 13
[INFO] [stdout]   %v6.v14 = insertelement <16 x i8> %v6.v13, i8 %v1, i32 14
[INFO] [stdout]   %v6 = insertelement <16 x i8> %v6.v14, i8 %v1, i32 15
[INFO] [stdout]   %v7.v0 = insertelement <8 x i8> poison, i8 %v3, i32 0
[INFO] [stdout]   %v7.v1 = insertelement <8 x i8> %v7.v0, i8 %v3, i32 1
[INFO] [stdout]   %v7.v2 = insertelement <8 x i8> %v7.v1, i8 %v3, i32 2
[INFO] [stdout]   %v7.v3 = insertelement <8 x i8> %v7.v2, i8 %v3, i32 3
[INFO] [stdout]   %v7.v4 = insertelement <8 x i8> %v7.v3, i8 %v3, i32 4
[INFO] [stdout]   %v7.v5 = insertelement <8 x i8> %v7.v4, i8 %v3, i32 5
[INFO] [stdout]   %v7.v6 = insertelement <8 x i8> %v7.v5, i8 %v3, i32 6
[INFO] [stdout]   %v7 = insertelement <8 x i8> %v7.v6, i8 %v3, i32 7
[INFO] [stdout]   %v8.v0 = insertelement <4 x float> poison, float %v5, i32 0
[INFO] [stdout]   %v8.v1 = insertelement <4 x float> %v8.v0, float %v5, i32 1
[INFO] [stdout]   %v8.v2 = insertelement <4 x float> %v8.v1, float %v5, i32 2
[INFO] [stdout]   %v8 = insertelement <4 x float> %v8.v2, float %v5, i32 3
[INFO] [stdout]   %v9 = getelementptr inbounds float, ptr %data3, i32 0
[INFO] [stdout]   %v10.aw = bitcast <16 x i8> %v6 to <4 x i32>
[INFO] [stdout]   %v10.a0 = extractelement <4 x i32> %v10.aw, i32 0
[INFO] [stdout]   %v10.a1 = extractelement <4 x i32> %v10.aw, i32 1
[INFO] [stdout]   %v10.a2 = extractelement <4 x i32> %v10.aw, i32 2
[INFO] [stdout]   %v10.a3 = extractelement <4 x i32> %v10.aw, i32 3
[INFO] [stdout]   %v10.bw = bitcast <8 x i8> %v7 to <2 x i32>
[INFO] [stdout]   %v10.b0 = extractelement <2 x i32> %v10.bw, i32 0
[INFO] [stdout]   %v10.b1 = extractelement <2 x i32> %v10.bw, i32 1
[INFO] [stdout]   %v10.c0 = extractelement <4 x float> %v8, i32 0
[INFO] [stdout]   %v10.c1 = extractelement <4 x float> %v8, i32 1
[INFO] [stdout]   %v10.c2 = extractelement <4 x float> %v8, i32 2
[INFO] [stdout]   %v10.c3 = extractelement <4 x float> %v8, i32 3
[INFO] [stdout]   %v10.mma = call { float, float, float, float } @llvm.nvvm.mma.m16n8k32.row.col.f32.e5m2.e5m2.f32(i32 %v10.a0, i32 %v10.a1, i32 %v10.a2, i32 %v10.a3, i32 %v10.b0, i32 %v10.b1, float %v10.c0, float %v10.c1, float %v10.c2, float %v10.c3)
[INFO] [stdout]   %v10.d0 = extractvalue { float, float, float, float } %v10.mma, 0
[INFO] [stdout]   %v10.d1 = extractvalue { float, float, float, float } %v10.mma, 1
[INFO] [stdout]   %v10.d2 = extractvalue { float, float, float, float } %v10.mma, 2
[INFO] [stdout]   %v10.d3 = extractvalue { float, float, float, float } %v10.mma, 3
[INFO] [stdout]   %v10.i0 = insertelement <4 x float> poison, float %v10.d0, i32 0
[INFO] [stdout]   %v10.i1 = insertelement <4 x float> %v10.i0, float %v10.d1, i32 1
[INFO] [stdout]   %v10.i2 = insertelement <4 x float> %v10.i1, float %v10.d2, i32 2
[INFO] [stdout]   %v10 = insertelement <4 x float> %v10.i2, float %v10.d3, i32 3
[INFO] [stdout]   %v11 = extractelement <4 x float> %v10, i32 0
[INFO] [stdout]   %v12 = extractelement <4 x float> %v10, i32 1
[INFO] [stdout]   %v13 = extractelement <4 x float> %v10, i32 2
[INFO] [stdout]   %v14 = extractelement <4 x float> %v10, i32 3
[INFO] [stdout]   store float %v11, ptr %v9
[INFO] [stdout]   store float %v12, ptr %v9
[INFO] [stdout]   store float %v13, ptr %v9
[INFO] [stdout]   store float %v14, ptr %v9
[INFO] [stdout]   ret void
[INFO] [stdout] }
[INFO] [stdout] 
[INFO] [stdout] attributes #0 = { nounwind "no-builtins" "no-trapping-math"="true" "nvvm.maxntid"="1" }
[INFO] [stdout] 
[INFO] [stdout] stack backtrace:
[INFO] [stdout]    0:     0x56754b60b151 - <<std[617be2ca0a6fd877]::sys::backtrace::BacktraceLock>::print::DisplayBacktrace as core[a373ad73a504971]::fmt::Display>::fmt
[INFO] [stdout]    1:     0x56754b621cca - core[a373ad73a504971]::fmt::write
[INFO] [stdout]    2:     0x56754b60f4ac - <alloc[2e1b2a31b91974c5]::vec::Vec<u8> as core[a373ad73a504971]::io::write::Write>::write_fmt
[INFO] [stdout]    3:     0x56754b5e60b6 - std[617be2ca0a6fd877]::panicking::default_hook::{closure#0}
[INFO] [stdout]    4:     0x56754b603379 - std[617be2ca0a6fd877]::panicking::default_hook
[INFO] [stdout]    5:     0x56754b31ddb0 - test[aada8e7e76ba119b]::test_main_inner::{closure#0}
[INFO] [stdout]    6:     0x56754b603532 - std[617be2ca0a6fd877]::panicking::panic_with_hook
[INFO] [stdout]    7:     0x56754b5e6162 - std[617be2ca0a6fd877]::panicking::panic_handler::{closure#0}
[INFO] [stdout]    8:     0x56754b5ded59 - std[617be2ca0a6fd877]::sys::backtrace::__rust_end_short_backtrace::<std[617be2ca0a6fd877]::panicking::panic_handler::{closure#0}, !>
[INFO] [stdout]    9:     0x56754b5e6d6d - __rustc[ea4f61d10b84890b]::rust_begin_unwind
[INFO] [stdout]   10:     0x56754b62242c - core[a373ad73a504971]::panicking::panic_fmt
[INFO] [stdout]   11:     0x56754b30487f - svod_codegen[ee8b13811706b789]::llvm::nvptx::ops::tests::assert_ptx_compiles
[INFO] [stdout]   12:     0x56754b2bb172 - svod_codegen[ee8b13811706b789]::llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing
[INFO] [stdout]   13:     0x56754b310e20 - <svod_codegen[ee8b13811706b789]::llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing::test_case_8_cuda_81632_e5m2_on_ada::{closure#0} as core[a373ad73a504971]::ops::function::FnOnce<()>>::call_once
[INFO] [stdout]   14:     0x56754b3111fb - test[aada8e7e76ba119b]::__rust_begin_short_backtrace::<core[a373ad73a504971]::result::Result<(), alloc[2e1b2a31b91974c5]::string::String>, fn() -> core[a373ad73a504971]::result::Result<(), alloc[2e1b2a31b91974c5]::string::String>>
[INFO] [stdout]   15:     0x56754b31e3d5 - test[aada8e7e76ba119b]::run_test::{closure#0}
[INFO] [stdout]   16:     0x56754b317cc4 - std[617be2ca0a6fd877]::sys::backtrace::__rust_begin_short_backtrace::<test[aada8e7e76ba119b]::run_test::{closure#1}, ()>
[INFO] [stdout]   17:     0x56754b321462 - <std[617be2ca0a6fd877]::thread::lifecycle::spawn_unchecked<test[aada8e7e76ba119b]::run_test::{closure#1}, ()>::{closure#1} as core[a373ad73a504971]::ops::function::FnOnce<()>>::call_once::{shim:vtable#0}
[INFO] [stdout]   18:     0x56754b60a79f - <std[617be2ca0a6fd877]::sys::thread::unix::Thread>::new::thread_start
[INFO] [stdout]   19:     0x78d49813adfa - <unknown>
[INFO] [stdout]   20:     0x78d4981ce3d4 - clone
[INFO] [stdout]   21:                0x0 - <unknown>
[INFO] [stdout] 
[INFO] [stdout] ---- llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing::test_case_7_cuda_81632_e4m3_on_ada stdout ----
[INFO] [stdout] 
[INFO] [stdout] thread 'llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing::test_case_7_cuda_81632_e4m3_on_ada' (2539) panicked at src/llvm/nvptx/../../test/unit/llvm_nvptx_ops.rs:76:5:
[INFO] [stdout] PTX carries an unresolved intrinsic call:
[INFO] [stdout] //
[INFO] [stdout] // Generated by LLVM NVPTX Back-End
[INFO] [stdout] //
[INFO] [stdout] 
[INFO] [stdout] .version 8.4
[INFO] [stdout] .target sm_89
[INFO] [stdout] .address_size 64
[INFO] [stdout] 
[INFO] [stdout] 	// .globl	nvptx_wmma              // -- Begin function nvptx_wmma
[INFO] [stdout] .extern .func  (.param .align 4 .b8 func_retval0[16]) llvm.nvvm.mma.m16n8k32.row.col.f32.e4m3.e4m3.f32
[INFO] [stdout] (
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e4m3.e4m3.f32_param_0,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e4m3.e4m3.f32_param_1,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e4m3.e4m3.f32_param_2,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e4m3.e4m3.f32_param_3,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e4m3.e4m3.f32_param_4,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e4m3.e4m3.f32_param_5,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e4m3.e4m3.f32_param_6,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e4m3.e4m3.f32_param_7,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e4m3.e4m3.f32_param_8,
[INFO] [stdout] 	.param .b32 llvm.nvvm.mma.m16n8k32.row.col.f32.e4m3.e4m3.f32_param_9
[INFO] [stdout] )
[INFO] [stdout] ;
[INFO] [stdout]                                         // @nvptx_wmma
[INFO] [stdout] .visible .entry nvptx_wmma(
[INFO] [stdout] 	.param .u64 .ptr .align 32 nvptx_wmma_param_0,
[INFO] [stdout] 	.param .u64 .ptr .align 32 nvptx_wmma_param_1,
[INFO] [stdout] 	.param .u64 .ptr .align 32 nvptx_wmma_param_2,
[INFO] [stdout] 	.param .u64 .ptr .align 32 nvptx_wmma_param_3
[INFO] [stdout] )
[INFO] [stdout] .maxntid 1
[INFO] [stdout] {
[INFO] [stdout] 	.reg .b32 	%r<13>;
[INFO] [stdout] 	.reg .b64 	%rd<9>;
[INFO] [stdout] 
[INFO] [stdout] // %bb.0:                               // %entry
[INFO] [stdout] 	ld.param.b64 	%rd1, [nvptx_wmma_param_0];
[INFO] [stdout] 	cvta.to.global.u64 	%rd2, %rd1;
[INFO] [stdout] 	ld.param.b64 	%rd3, [nvptx_wmma_param_1];
[INFO] [stdout] 	cvta.to.global.u64 	%rd4, %rd3;
[INFO] [stdout] 	ld.param.b64 	%rd5, [nvptx_wmma_param_2];
[INFO] [stdout] 	cvta.to.global.u64 	%rd6, %rd5;
[INFO] [stdout] 	ld.param.b64 	%rd7, [nvptx_wmma_param_3];
[INFO] [stdout] 	cvta.to.global.u64 	%rd8, %rd7;
[INFO] [stdout] 	ld.global.nc.b32 	%r1, [%rd6];
[INFO] [stdout] 	ld.global.nc.b8 	%r2, [%rd2];
[INFO] [stdout] 	prmt.b32 	%r3, %r2, %r2, 0x3340U;
[INFO] [stdout] 	prmt.b32 	%r4, %r3, %r3, 0x5410U;
[INFO] [stdout] 	ld.global.nc.b8 	%r5, [%rd4];
[INFO] [stdout] 	prmt.b32 	%r6, %r5, %r5, 0x3340U;
[INFO] [stdout] 	prmt.b32 	%r7, %r6, %r6, 0x5410U;
[INFO] [stdout] 	{ // callseq 0, 0
[INFO] [stdout] 	.param .b32 	param0;
[INFO] [stdout] 	st.param.b32 	[param0], %r4;
[INFO] [stdout] 	.param .b32 	param1;
[INFO] [stdout] 	st.param.b32 	[param1], %r4;
[INFO] [stdout] 	.param .b32 	param2;
[INFO] [stdout] 	st.param.b32 	[param2], %r4;
[INFO] [stdout] 	.param .b32 	param3;
[INFO] [stdout] 	st.param.b32 	[param3], %r4;
[INFO] [stdout] 	.param .b32 	param4;
[INFO] [stdout] 	st.param.b32 	[param4], %r7;
[INFO] [stdout] 	.param .b32 	param5;
[INFO] [stdout] 	st.param.b32 	[param5], %r7;
[INFO] [stdout] 	.param .b32 	param6;
[INFO] [stdout] 	st.param.b32 	[param6], %r1;
[INFO] [stdout] 	.param .b32 	param7;
[INFO] [stdout] 	st.param.b32 	[param7], %r1;
[INFO] [stdout] 	.param .b32 	param8;
[INFO] [stdout] 	st.param.b32 	[param8], %r1;
[INFO] [stdout] 	.param .b32 	param9;
[INFO] [stdout] 	st.param.b32 	[param9], %r1;
[INFO] [stdout] 	.param .align 4 .b8 	retval0[16];
[INFO] [stdout] 	call.uni (retval0), llvm.nvvm.mma.m16n8k32.row.col.f32.e4m3.e4m3.f32, (param0, param1, param2, param3, param4, param5, param6, param7, param8, param9);
[INFO] [stdout] 	ld.param.b32 	%r8, [retval0];
[INFO] [stdout] 	ld.param.b32 	%r9, [retval0+4];
[INFO] [stdout] 	ld.param.b32 	%r10, [retval0+8];
[INFO] [stdout] 	ld.param.b32 	%r11, [retval0+12];
[INFO] [stdout] 	} // callseq 0
[INFO] [stdout] 	st.global.b32 	[%rd8], %r11;
[INFO] [stdout] 	ret;
[INFO] [stdout]                                         // -- End function
[INFO] [stdout] }
[INFO] [stdout] 
[INFO] [stdout] --- IR ---
[INFO] [stdout] ; ModuleID = 'nvptx_wmma'
[INFO] [stdout] source_filename = "nvptx_wmma"
[INFO] [stdout] target datalayout = "e-p6:32:32-i64:64-i128:128-i256:256-v16:16-v32:32-n16:32:64"
[INFO] [stdout] target triple = "nvptx64-nvidia-cuda"
[INFO] [stdout] 
[INFO] [stdout] declare { float, float, float, float } @llvm.nvvm.mma.m16n8k32.row.col.f32.e4m3.e4m3.f32(i32, i32, i32, i32, i32, i32, float, float, float, float)
[INFO] [stdout] 
[INFO] [stdout] define ptx_kernel void @nvptx_wmma(ptr noalias align 32 %data0, ptr noalias align 32 %data1, ptr noalias align 32 %data2, ptr noalias align 32 %data3) #0 {
[INFO] [stdout] entry:
[INFO] [stdout] 
[INFO] [stdout]   %v0 = getelementptr inbounds i8, ptr %data0, i32 0
[INFO] [stdout]   %v1 = load i8, ptr %v0
[INFO] [stdout]   %v2 = getelementptr inbounds i8, ptr %data1, i32 0
[INFO] [stdout]   %v3 = load i8, ptr %v2
[INFO] [stdout]   %v4 = getelementptr inbounds float, ptr %data2, i32 0
[INFO] [stdout]   %v5 = load float, ptr %v4
[INFO] [stdout]   %v6.v0 = insertelement <16 x i8> poison, i8 %v1, i32 0
[INFO] [stdout]   %v6.v1 = insertelement <16 x i8> %v6.v0, i8 %v1, i32 1
[INFO] [stdout]   %v6.v2 = insertelement <16 x i8> %v6.v1, i8 %v1, i32 2
[INFO] [stdout]   %v6.v3 = insertelement <16 x i8> %v6.v2, i8 %v1, i32 3
[INFO] [stdout]   %v6.v4 = insertelement <16 x i8> %v6.v3, i8 %v1, i32 4
[INFO] [stdout]   %v6.v5 = insertelement <16 x i8> %v6.v4, i8 %v1, i32 5
[INFO] [stdout]   %v6.v6 = insertelement <16 x i8> %v6.v5, i8 %v1, i32 6
[INFO] [stdout]   %v6.v7 = insertelement <16 x i8> %v6.v6, i8 %v1, i32 7
[INFO] [stdout]   %v6.v8 = insertelement <16 x i8> %v6.v7, i8 %v1, i32 8
[INFO] [stdout]   %v6.v9 = insertelement <16 x i8> %v6.v8, i8 %v1, i32 9
[INFO] [stdout]   %v6.v10 = insertelement <16 x i8> %v6.v9, i8 %v1, i32 10
[INFO] [stdout]   %v6.v11 = insertelement <16 x i8> %v6.v10, i8 %v1, i32 11
[INFO] [stdout]   %v6.v12 = insertelement <16 x i8> %v6.v11, i8 %v1, i32 12
[INFO] [stdout]   %v6.v13 = insertelement <16 x i8> %v6.v12, i8 %v1, i32 13
[INFO] [stdout]   %v6.v14 = insertelement <16 x i8> %v6.v13, i8 %v1, i32 14
[INFO] [stdout]   %v6 = insertelement <16 x i8> %v6.v14, i8 %v1, i32 15
[INFO] [stdout]   %v7.v0 = insertelement <8 x i8> poison, i8 %v3, i32 0
[INFO] [stdout]   %v7.v1 = insertelement <8 x i8> %v7.v0, i8 %v3, i32 1
[INFO] [stdout]   %v7.v2 = insertelement <8 x i8> %v7.v1, i8 %v3, i32 2
[INFO] [stdout]   %v7.v3 = insertelement <8 x i8> %v7.v2, i8 %v3, i32 3
[INFO] [stdout]   %v7.v4 = insertelement <8 x i8> %v7.v3, i8 %v3, i32 4
[INFO] [stdout]   %v7.v5 = insertelement <8 x i8> %v7.v4, i8 %v3, i32 5
[INFO] [stdout]   %v7.v6 = insertelement <8 x i8> %v7.v5, i8 %v3, i32 6
[INFO] [stdout]   %v7 = insertelement <8 x i8> %v7.v6, i8 %v3, i32 7
[INFO] [stdout]   %v8.v0 = insertelement <4 x float> poison, float %v5, i32 0
[INFO] [stdout]   %v8.v1 = insertelement <4 x float> %v8.v0, float %v5, i32 1
[INFO] [stdout]   %v8.v2 = insertelement <4 x float> %v8.v1, float %v5, i32 2
[INFO] [stdout]   %v8 = insertelement <4 x float> %v8.v2, float %v5, i32 3
[INFO] [stdout]   %v9 = getelementptr inbounds float, ptr %data3, i32 0
[INFO] [stdout]   %v10.aw = bitcast <16 x i8> %v6 to <4 x i32>
[INFO] [stdout]   %v10.a0 = extractelement <4 x i32> %v10.aw, i32 0
[INFO] [stdout]   %v10.a1 = extractelement <4 x i32> %v10.aw, i32 1
[INFO] [stdout]   %v10.a2 = extractelement <4 x i32> %v10.aw, i32 2
[INFO] [stdout]   %v10.a3 = extractelement <4 x i32> %v10.aw, i32 3
[INFO] [stdout]   %v10.bw = bitcast <8 x i8> %v7 to <2 x i32>
[INFO] [stdout]   %v10.b0 = extractelement <2 x i32> %v10.bw, i32 0
[INFO] [stdout]   %v10.b1 = extractelement <2 x i32> %v10.bw, i32 1
[INFO] [stdout]   %v10.c0 = extractelement <4 x float> %v8, i32 0
[INFO] [stdout]   %v10.c1 = extractelement <4 x float> %v8, i32 1
[INFO] [stdout]   %v10.c2 = extractelement <4 x float> %v8, i32 2
[INFO] [stdout]   %v10.c3 = extractelement <4 x float> %v8, i32 3
[INFO] [stdout]   %v10.mma = call { float, float, float, float } @llvm.nvvm.mma.m16n8k32.row.col.f32.e4m3.e4m3.f32(i32 %v10.a0, i32 %v10.a1, i32 %v10.a2, i32 %v10.a3, i32 %v10.b0, i32 %v10.b1, float %v10.c0, float %v10.c1, float %v10.c2, float %v10.c3)
[INFO] [stdout]   %v10.d0 = extractvalue { float, float, float, float } %v10.mma, 0
[INFO] [stdout]   %v10.d1 = extractvalue { float, float, float, float } %v10.mma, 1
[INFO] [stdout]   %v10.d2 = extractvalue { float, float, float, float } %v10.mma, 2
[INFO] [stdout]   %v10.d3 = extractvalue { float, float, float, float } %v10.mma, 3
[INFO] [stdout]   %v10.i0 = insertelement <4 x float> poison, float %v10.d0, i32 0
[INFO] [stdout]   %v10.i1 = insertelement <4 x float> %v10.i0, float %v10.d1, i32 1
[INFO] [stdout]   %v10.i2 = insertelement <4 x float> %v10.i1, float %v10.d2, i32 2
[INFO] [stdout]   %v10 = insertelement <4 x float> %v10.i2, float %v10.d3, i32 3
[INFO] [stdout]   %v11 = extractelement <4 x float> %v10, i32 0
[INFO] [stdout]   %v12 = extractelement <4 x float> %v10, i32 1
[INFO] [stdout]   %v13 = extractelement <4 x float> %v10, i32 2
[INFO] [stdout]   %v14 = extractelement <4 x float> %v10, i32 3
[INFO] [stdout]   store float %v11, ptr %v9
[INFO] [stdout]   store float %v12, ptr %v9
[INFO] [stdout]   store float %v13, ptr %v9
[INFO] [stdout]   store float %v14, ptr %v9
[INFO] [stdout]   ret void
[INFO] [stdout] }
[INFO] [stdout] 
[INFO] [stdout] attributes #0 = { nounwind "no-builtins" "no-trapping-math"="true" "nvvm.maxntid"="1" }
[INFO] [stdout] 
[INFO] [stdout] stack backtrace:
[INFO] [stdout]    0:     0x56754b60b151 - <<std[617be2ca0a6fd877]::sys::backtrace::BacktraceLock>::print::DisplayBacktrace as core[a373ad73a504971]::fmt::Display>::fmt
[INFO] [stdout]    1:     0x56754b621cca - core[a373ad73a504971]::fmt::write
[INFO] [stdout]    2:     0x56754b60f4ac - <alloc[2e1b2a31b91974c5]::vec::Vec<u8> as core[a373ad73a504971]::io::write::Write>::write_fmt
[INFO] [stdout]    3:     0x56754b5e60b6 - std[617be2ca0a6fd877]::panicking::default_hook::{closure#0}
[INFO] [stdout]    4:     0x56754b603379 - std[617be2ca0a6fd877]::panicking::default_hook
[INFO] [stdout]    5:     0x56754b31ddb0 - test[aada8e7e76ba119b]::test_main_inner::{closure#0}
[INFO] [stdout]    6:     0x56754b603532 - std[617be2ca0a6fd877]::panicking::panic_with_hook
[INFO] [stdout]    7:     0x56754b5e6162 - std[617be2ca0a6fd877]::panicking::panic_handler::{closure#0}
[INFO] [stdout]    8:     0x56754b5ded59 - std[617be2ca0a6fd877]::sys::backtrace::__rust_end_short_backtrace::<std[617be2ca0a6fd877]::panicking::panic_handler::{closure#0}, !>
[INFO] [stdout]    9:     0x56754b5e6d6d - __rustc[ea4f61d10b84890b]::rust_begin_unwind
[INFO] [stdout]   10:     0x56754b62242c - core[a373ad73a504971]::panicking::panic_fmt
[INFO] [stdout]   11:     0x56754b30487f - svod_codegen[ee8b13811706b789]::llvm::nvptx::ops::tests::assert_ptx_compiles
[INFO] [stdout]   12:     0x56754b2bb172 - svod_codegen[ee8b13811706b789]::llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing
[INFO] [stdout]   13:     0x56754b310d60 - <svod_codegen[ee8b13811706b789]::llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing::test_case_7_cuda_81632_e4m3_on_ada::{closure#0} as core[a373ad73a504971]::ops::function::FnOnce<()>>::call_once
[INFO] [stdout]   14:     0x56754b3111fb - test[aada8e7e76ba119b]::__rust_begin_short_backtrace::<core[a373ad73a504971]::result::Result<(), alloc[2e1b2a31b91974c5]::string::String>, fn() -> core[a373ad73a504971]::result::Result<(), alloc[2e1b2a31b91974c5]::string::String>>
[INFO] [stdout]   15:     0x56754b31e3d5 - test[aada8e7e76ba119b]::run_test::{closure#0}
[INFO] [stdout]   16:     0x56754b317cc4 - std[617be2ca0a6fd877]::sys::backtrace::__rust_begin_short_backtrace::<test[aada8e7e76ba119b]::run_test::{closure#1}, ()>
[INFO] [stdout]   17:     0x56754b321462 - <std[617be2ca0a6fd877]::thread::lifecycle::spawn_unchecked<test[aada8e7e76ba119b]::run_test::{closure#1}, ()>::{closure#1} as core[a373ad73a504971]::ops::function::FnOnce<()>>::call_once::{shim:vtable#0}
[INFO] [stdout]   18:     0x56754b60a79f - <std[617be2ca0a6fd877]::sys::thread::unix::Thread>::new::thread_start
[INFO] [stdout]   19:     0x78d49813adfa - <unknown>
[INFO] [stdout]   20:     0x78d4981ce3d4 - clone
[INFO] [stdout]   21:                0x0 - <unknown>
[INFO] [stdout] 
[INFO] [stdout] 
[INFO] [stdout] failures:
[INFO] [stdout]     llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing::test_case_7_cuda_81632_e4m3_on_ada
[INFO] [stdout]     llvm::nvptx::wmma::tests::nvptx_wmma_operand_packing::test_case_8_cuda_81632_e5m2_on_ada
[INFO] [stdout] 
[INFO] [stdout] test result: FAILED. 343 passed; 2 failed; 0 ignored; 0 measured; 0 filtered out; finished in 1.93s
[INFO] [stdout] 
[INFO] running `Command { std: "docker" "inspect" "2c77b06722d8a908026bfcb438b97adc842c6fc92fbd91244fc7d443de04deac", kill_on_drop: false }`
[INFO] running `Command { std: "docker" "rm" "-f" "2c77b06722d8a908026bfcb438b97adc842c6fc92fbd91244fc7d443de04deac", kill_on_drop: false }`
[INFO] [stdout] 2c77b06722d8a908026bfcb438b97adc842c6fc92fbd91244fc7d443de04deac
