// -----// IR Dump Before ConvertToNVVMPass: iree-convert-to-nvvm //----- //
module {
  func.func @main_graph$async_dispatch_1_elementwise_D_i64xi1() attributes {gpu.known_block_size = array<i32: 32, 1, 1>} {
    %c32_i32 = arith.constant 32 : i32
    %c-32 = arith.constant -32 : index
    %c1 = arith.constant 1 : index
    %c32 = arith.constant 32 : index
    %c0 = arith.constant 0 : index
    %c32_i64 = arith.constant 32 : i64
    %cst = arith.constant dense<0> : vector<1xi64>
    %thread_id_x = gpu.thread_id x upper_bound 32
    %0 = hal.interface.constant.load layout(<constants = 2, bindings = [#hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, Indirect>], flags = Indirect>) ordinal(0) : i32
    %1 = hal.interface.constant.load layout(<constants = 2, bindings = [#hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, Indirect>], flags = Indirect>) ordinal(1) : i32
    %2 = arith.extui %0 : i32 to i64
    %3 = arith.extui %1 : i32 to i64
    %4 = arith.shli %3, %c32_i64 : i64
    %5 = arith.ori %2, %4 : i64
    %6 = arith.index_castui %5 : i64 to index
    %7 = util.assume.int %6<umin = 0, umax = 9223372036854775807> : index
    %8 = hal.interface.binding.subspan layout(<constants = 2, bindings = [#hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, Indirect>], flags = Indirect>) binding(0) alignment(64) offset(%c0) flags("ReadOnly|Indirect") : memref<?xi64, #gpu.address_space<global>>{%7}
    %assume_align = memref.assume_alignment %8, 64 : memref<?xi64, #gpu.address_space<global>>
    %9 = hal.interface.binding.subspan layout(<constants = 2, bindings = [#hal.pipeline.binding<storage_buffer, "ReadOnly|Indirect">, #hal.pipeline.binding<storage_buffer, Indirect>], flags = Indirect>) binding(1) alignment(64) offset(%c0) flags(Indirect) : memref<?xi8, #gpu.address_space<global>>{%7}
    %assume_align_0 = memref.assume_alignment %9, 64 : memref<?xi8, #gpu.address_space<global>>
    %10 = arith.cmpi sle, %7, %c0 : index
    %11 = arith.subi %c0, %7 : index
    %12 = arith.subi %7, %c1 : index
    %13 = arith.select %10, %11, %12 : index
    %14 = arith.divsi %13, %c32 : index
    %15 = arith.subi %c0, %14 : index
    %16 = arith.addi %14, %c1 : index
    %17 = arith.select %10, %15, %16 : index
    %workgroup_id_x = hal.interface.workgroup.id[0] upper_bound 2147483647 : index
    %workgroup_count_x = hal.interface.workgroup.count[0] upper_bound 2147483647 : index
    cf.br ^bb1(%workgroup_id_x : index)
  ^bb1(%18: index):  // 2 preds: ^bb0, ^bb5
    %19 = arith.cmpi slt, %18, %17 : index
    cf.cond_br %19, ^bb2, ^bb6
  ^bb2:  // pred: ^bb1
    %20 = arith.muli %18, %c-32 overflow<nsw> : index
    %21 = arith.addi %20, %7 : index
    %22 = arith.minsi %21, %c32 : index
    %23 = arith.index_castui %thread_id_x : index to i32
    %24 = arith.index_castui %22 : index to i32
    cf.br ^bb3(%23 : i32)
  ^bb3(%25: i32):  // 2 preds: ^bb2, ^bb4
    %26 = arith.cmpi slt, %25, %24 : i32
    cf.cond_br %26, ^bb4, ^bb5
  ^bb4:  // pred: ^bb3
    %27 = arith.index_castui %25 : i32 to index
    %28 = arith.muli %18, %c32 overflow<nsw> : index
    %29 = arith.addi %28, %27 : index
    %30 = vector.load %assume_align[%29] : memref<?xi64, #gpu.address_space<global>>, vector<1xi64>
    %31 = arith.cmpi ne, %30, %cst : vector<1xi64>
    %32 = arith.extui %31 : vector<1xi1> to vector<1xi8>
    vector.store %32, %assume_align_0[%29] : memref<?xi8, #gpu.address_space<global>>, vector<1xi8>
    %33 = arith.addi %25, %c32_i32 : i32
    cf.br ^bb3(%33 : i32)
  ^bb5:  // pred: ^bb3
    %34 = arith.addi %18, %workgroup_count_x : index
    cf.br ^bb1(%34 : index)
  ^bb6:  // pred: ^bb1
    return
  }
  iree_codegen.dispatch_config @main_graph$async_dispatch_1_elementwise_D_i64xi1 workgroup_size = [32, 1, 1] subgroup_size = 32 {
  ^bb0(%arg0: !hal.device, %arg1: index):
    %c1 = arith.constant 1 : index
    %0 = affine.min affine_map<()[s0] -> (s0 ceildiv 32, 2147483647)>()[%arg1]
    iree_codegen.yield %0, %c1, %c1 : index, index, index
  }
}