Bump to llvm/llvm-project@b959532

[MaheshRavishankar]

This PR contains revert of the following commits from upstream

• llvm/llvm-project@b6603e1 : Causes build failure in https://github.com/llvm/torch-mlir/blob/main/lib/Dialect/Torch/Transforms/InlineGlobalSlots.cpp
• iree-org/llvm-project@d7073c5 : Causes various errors with EmitC and failure in https://github.com/iree-org/iree/blob/main/compiler/src/iree/compiler/Codegen/Common/TypePropagationPass.cpp
• llvm/llvm-project@1387ba4 : Causes failure in https://github.com/iree-org/iree/blob/main/compiler/src/iree/compiler/Codegen/LLVMGPU/test/rocdl_pipeline_test.mlir

[kuhar]

ROCm changes LGTM

[MaheshRavishankar]

The failure with llvm/llvm-project@1387ba4 can be reproed by this test

module {
  func.func @ext_fp8_dispatch() attributes {translation_info = #iree_codegen.translation_info<LLVMGPUVectorize workgroup_size = [128, 1, 1] subgroup_size = 64>} {
    %c128 = arith.constant 128 : index
    %c0 = arith.constant 0 : index
    %0 = hal.interface.binding.subspan layout(<bindings = [#hal.pipeline.binding<storage_buffer>, #hal.pipeline.binding<storage_buffer>, #hal.pipeline.binding<storage_buffer>]>) binding(0) alignment(64) offset(%c0) flags(ReadOnly) : me\mref<4096xf8E4M3FNUZ, #gpu.address_space<global>>
    memref.assume_alignment %0, 64 : memref<4096xf8E4M3FNUZ, #gpu.address_space<global>>
    %1 = hal.interface.binding.subspan layout(<bindings = [#hal.pipeline.binding<storage_buffer>, #hal.pipeline.binding<storage_buffer>, #hal.pipeline.binding<storage_buffer>]>) binding(1) alignment(64) offset(%c0) flags(ReadOnly) : me\mref<4096xf8E5M2FNUZ, #gpu.address_space<global>>
    memref.assume_alignment %1, 64 : memref<4096xf8E5M2FNUZ, #gpu.address_space<global>>
    %2 = hal.interface.binding.subspan layout(<bindings = [#hal.pipeline.binding<storage_buffer>, #hal.pipeline.binding<storage_buffer>, #hal.pipeline.binding<storage_buffer>]>) binding(2) alignment(64) offset(%c0) : memref<4096xf32, #\gpu.address_space<global>>
    memref.assume_alignment %2, 64 : memref<4096xf32, #gpu.address_space<global>>
    %workgroup_id_x = hal.interface.workgroup.id[0] : index
    %thread_id_x = gpu.thread_id  x
    %3 = arith.muli %workgroup_id_x, %c128 : index
    %4 = arith.addi %thread_id_x, %3 : index
    %5 = memref.load %0[%4] : memref<4096xf8E4M3FNUZ, #gpu.address_space<global>>
    %6 = memref.load %1[%4] : memref<4096xf8E5M2FNUZ, #gpu.address_space<global>>
    %7 = arith.extf %5 : f8E4M3FNUZ to f32
    %8 = arith.extf %6 : f8E5M2FNUZ to f32
    %9 = arith.addf %7, %8 : f32
    memref.store %9, %2[%4] : memref<4096xf32, #gpu.address_space<global>>
    return
  }
}

iree-opt  --iree-gpu-test-target=gfx908 --iree-convert-to-rocdl small_repro.mlir

[krzysz00]

@MaheshRavishankar I'd argue that that test should fail on gfx908 - gfx908 doesn't support fp8 types, which means that those extfs shouldn't lower.

I'm confused by why that ever passed and what the generated LLVM IR looked like

[MaheshRavishankar]

There might be some legalization in LLVM for these types on non-supported hardware. Would be good to check.