Gather/scatter emitted by gb main ci

[wsmoses]

    %c_108 = stablehlo.constant dense<16> : tensor<256xi64> loc(#loc645)

    %c_110 = stablehlo.constant dense<"0xFFFFFFFFFFFFFFFF0100000000000000020000000000000003000000000000000400000000000000050000000000000006000000000000000700000000000000080000000000000009000000000000000A000000000000000B000000000000000C000000000000000D000000000000000E000000000000000F0000000000000010000000000000001100000000000000120000000000000013000000000000001400000000000000150000000000000016000000000000001700000000000000180000000000000019000000000000001A000000000000001B000000000000001C000000000000001D000000000000001E000000000000001F0000000000000020000000000000002100000000000000220000000000000023000000000000002400000000000000250000000000000026000000000000002700000000000000280000000000000029000000000000002A000000000000002B000000000000002C000000000000002D000000000000002E000000000000002F0000000000000030000000000000003100000000000000320000000000000033000000000000003400000000000000350000000000000036000000000000003700000000000000380000000000000039000000000000003A000000000000003B000000000000003C000000000000003D000000000000003E000000000000003F0000000000000040000000000000004100000000000000420000000000000043000000000000004400000000000000450000000000000046000000000000004700000000000000480000000000000049000000000000004A000000000000004B000000000000004C000000000000004D000000000000004E000000000000004F0000000000000050000000000000005100000000000000520000000000000053000000000000005400000000000000550000000000000056000000000000005700000000000000580000000000000059000000000000005A000000000000005B000000000000005C000000000000005D000000000000005E000000000000005F0000000000000060000000000000006100000000000000620000000000000063000000000000006400000000000000650000000000000066000000000000006700000000000000680000000000000069000000000000006A000000000000006B000000000000006C000000000000006D000000000000006E000000000000006F0000000000000070000000000000007100000000000000720000000000000073000000000000007400000000000000750000000000000076000000000000007700000000000000780000000000000079000000000000007A000000000000007B000000000000007C000000000000007D000000000000007E000000000000007F0000000000000080000000000000008100000000000000820000000000000083000000000000008400000000000000850000000000000086000000000000008700000000000000880000000000000089000000000000008A000000000000008B000000000000008C000000000000008D000000000000008E000000000000008F0000000000000090000000000000009100000000000000920000000000000093000000000000009400000000000000950000000000000096000000000000009700000000000000980000000000000099000000000000009A000000000000009B000000000000009C000000000000009D000000000000009E000000000000009F00000000000000A000000000000000A100000000000000A200000000000000A300000000000000A400000000000000A500000000000000A600000000000000A700000000000000A800000000000000A900000000000000AA00000000000000AB00000000000000AC00000000000000AD00000000000000AE00000000000000AF00000000000000B000000000000000B100000000000000B200000000000000B300000000000000B400000000000000B500000000000000B600000000000000B700000000000000B800000000000000B900000000000000BA00000000000000BB00000000000000BC00000000000000BD00000000000000BE00000000000000BF00000000000000C000000000000000C100000000000000C200000000000000C300000000000000C400000000000000C500000000000000C600000000000000C700000000000000C800000000000000C900000000000000CA00000000000000CB00000000000000CC00000000000000CD00000000000000CE00000000000000CF00000000000000D000000000000000D100000000000000D200000000000000D300000000000000D400000000000000D500000000000000D600000000000000D700000000000000D800000000000000D900000000000000DA00000000000000DB00000000000000DC00000000000000DD00000000000000DE00000000000000DF00000000000000E000000000000000E100000000000000E200000000000000E300000000000000E400000000000000E500000000000000E600000000000000E700000000000000E800000000000000E900000000000000EA00000000000000EB00000000000000EC00000000000000ED00000000000000EE00000000000000EF00000000000000F000000000000000F100000000000000F200000000000000F300000000000000F400000000000000F500000000000000F600000000000000F700000000000000F800000000000000F900000000000000FA00000000000000FB00000000000000FC00000000000000FD00000000000000FE00000000000000FF00000000000000"> : tensor<256xi64> loc(#loc645)

    %c_88 = stablehlo.constant {enzymexla.non_negative = [#enzymexla<guaranteed GUARANTEED>]} dense<1> : tensor<256xi64> loc(#loc)

    %c_109 = stablehlo.constant dense<"0x01000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000000"> : tensor<256xi1> loc(#loc645)

    %c_85 = stablehlo.constant dense<[0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20, 21, 22, 23, 24, 25, 26, 27, 28, 29, 30, 31, 32, 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45, 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58, 59, 60, 61, 62, 63, 64, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80, 81, 82, 83, 84, 85, 86, 87, 88, 89, 90, 91, 92, 93, 94]> : tensor<95xi64> loc(#loc1135)

    %c_332 = stablehlo.constant dense<16> : tensor<95xi64> loc(#loc925)

    %25 = stablehlo.divide %c_110, %c_108 : tensor<256xi64> loc(#loc645)
    %26 = stablehlo.negate %25 : tensor<256xi64> loc(#loc645)

    %27 = stablehlo.subtract %26, %c_88 {enzymexla.non_negative = [#enzymexla<guaranteed NOTGUARANTEED>]} : tensor<256xi64> loc(#loc645)

    %28 = stablehlo.select %c_109, %27, %25 {enzymexla.non_negative = [#enzymexla<guaranteed NOTGUARANTEED>]} : tensor<256xi1>, tensor<256xi64> loc(#loc645)


    %3067 = stablehlo.multiply %c_85, %c_332 : tensor<95xi64> loc(#loc1162)

    %3068 = stablehlo.broadcast_in_dim %3067, dims = [0] : (tensor<95xi64>) -> tensor<95x256xi64> loc(#loc925)
    %3069 = stablehlo.broadcast_in_dim %28, dims = [1] : (tensor<256xi64>) -> tensor<95x256xi64> loc(#loc925)

    %3070 = stablehlo.add %3068, %3069 : tensor<95x256xi64> loc(#loc925)
    %3071 = stablehlo.reshape %3070 : (tensor<95x256xi64>) -> tensor<24320x1xi64> loc(#loc925)

    %3083 = stablehlo.add %3071, %c_334 : tensor<24320x1xi64> loc(#loc925)
    %3084 = "stablehlo.gather"(%arg6, %3083) <{dimension_numbers = #stablehlo.gather<collapsed_slice_dims = [0], start_index_map = [0], index_vector_dim = 1>, indices_are_sorted = false, slice_sizes = array<i64: 1>}> : (tensor<1536xf32>, tensor<24320x1xi64>) -> tensor<24320xf32> loc(#loc925)

cc @avik-pal @dkytezab . Bad regardless, may be part of the cause of current pain

[wsmoses]

possible partial fix [for this gather] in #2498, also scatter exists if you look at mlir logs

[wsmoses]

@dkytezab can you help find the input IR / reason from julia why now all of a sudden there are scatters/gathers coming out?

concurrently @avik-pal we should get the fixups in

I'm also confused why we are failing to constprop here effectively

[wsmoses]

also cc @giordano / @gbaraldi this may be a cuda.jl version thing again.

  llvm.func local_unnamed_addr @_Z29gpu__compute_barotropic_mode_16CompilerMetadataI10StaticSizeI12_3056__1520_E12DynamicCheckvv7NDRangeILi2ES0_I9_191__95_ES0_I8_16__16_EvvEE5FieldI4Face6Centervvvv11OffsetArrayI7Float32Li3E13CuTracedArrayISC_Li3ELi1E15_3072__1536__1_EESC_vvvES8_ISA_S9_vvvvSF_SC_vvvE21LatitudeLongitudeGridI15CuTracedRNumberISC_Li1EE8Periodic7BoundedSM_28StaticVerticalDiscretizationISB_ISC_Li1ESD_ISC_Li1ELi1E5_21__EESB_ISC_Li1ESD_ISC_Li1ELi1E5_20__EESP_SR_ESC_SC_SB_ISC_Li1E12StepRangeLenISC_7Float64SU_5Int64EESX_SC_SC_SX_SX_SB_ISC_Li1ESD_ISC_Li1ELi1E7_1536__EESZ_SZ_SZ_SC_SC_vSV_ESB_ISC_Li3ESD_ISC_Li3ELi1E16_3072__1536__20_EES12_SF_(%arg0: !llvm.struct<(struct<(array<1 x ptr<1>>, array<3 x i64>)>)> loc(fused<#di_subprogram>[#loc3]), %arg1: !llvm.struct<(struct<(array<1 x ptr<1>>, array<3 x i64>)>)> loc(fused<#di_subprogram>[#loc3]), %arg2: !llvm.struct<(i64, i64, i64, i64, i64, i64, array<1 x ptr<1>>, array<1 x ptr<1>>, array<1 x ptr<1>>, f32, f32, struct<(struct<(f64, f64, i64, i64)>, array<1 x i64>)>, struct<(struct<(f64, f64, i64, i64)>, array<1 x i64>)>, f32, f32, struct<(struct<(f64, f64, i64, i64)>, array<1 x i64>)>, struct<(struct<(f64, f64, i64, i64)>, array<1 x i64>)>, array<4 x struct<(array<1 x ptr<1>>, array<1 x i64>)>>, struct<(array<1 x ptr<1>>, array<1 x i64>)>, struct<(array<1 x ptr<1>>, array<1 x i64>)>, struct<(array<1 x ptr<1>>, array<1 x i64>)>, struct<(array<1 x ptr<1>>, array<1 x i64>)>, f32, f32, struct<(array<1 x ptr<1>>, array<1 x i64>)>, struct<(array<1 x ptr<1>>, array<1 x i64>)>, struct<(array<1 x ptr<1>>, array<1 x i64>)>, struct<(array<1 x ptr<1>>, array<1 x i64>)>, array<1 x ptr<1>>)> loc(fused<#di_subprogram>[#loc3]), %arg3: !llvm.struct<(array<1 x ptr<1>>, array<3 x i64>)> loc(fused<#di_subprogram>[#loc3]), %arg4: !llvm.struct<(array<1 x ptr<1>>, array<3 x i64>)> loc(fused<#di_subprogram>[#loc3]), %arg5: !llvm.struct<(array<1 x ptr<1>>, array<3 x i64>)> loc(fused<#di_subprogram>[#loc3])) attributes {sym_visibility = "private"} {
    %0 = llvm.mlir.constant(191 : i32) : i32 loc(#loc)
    %1 = llvm.mlir.constant(1152921504606846785 : i64) : i64 loc(#loc)
    %2 = llvm.mlir.constant(4 : i64) : i64 loc(#loc)
    %3 = llvm.mlir.constant(1 : i64) : i64 loc(#loc)
    %4 = llvm.mlir.constant(3055 : i64) : i64 loc(#loc)
    %5 = llvm.mlir.constant(1520 : i64) : i64 loc(#loc)
    %6 = llvm.mlir.constant(15 : i64) : i64 loc(#loc)
    %7 = llvm.mlir.constant(0.000000e+00 : f32) : f32 loc(#loc)
    %8 = llvm.mlir.constant(1.000000e+00 : f32) : f32 loc(#loc)
    %9 = llvm.mlir.constant(0 : i64) : i64 loc(#loc)
    %10 = llvm.mlir.constant(-1 : i64) : i64 loc(#loc)
    %11 = llvm.mlir.constant(3072 : i64) : i64 loc(#loc)
    %12 = llvm.mlir.constant(-4718592 : i64) : i64 loc(#loc)
    %13 = llvm.mlir.constant(2 : i64) : i64 loc(#loc)
    %14 = llvm.mlir.constant(4718592 : i64) : i64 loc(#loc)
    %15 = llvm.extractvalue %arg2[17, 3, 1, 0] : !llvm.struct<(i64, i64, i64, i64, i64, i64, array<1 x ptr<1>>, array<1 x ptr<1>>, array<1 x ptr<1>>, f32, f32, struct<(struct<(f64, f64, i64, i64)>, array<1 x i64>)>, struct<(struct<(f64, f64, i64, i64)>, array<1 x i64>)>, f32, f32, struct<(struct<(f64, f64, i64, i64)>, array<1 x i64>)>, struct<(struct<(f64, f64, i64, i64)>, array<1 x i64>)>, array<4 x struct<(array<1 x ptr<1>>, array<1 x i64>)>>, struct<(array<1 x ptr<1>>, array<1 x i64>)>, struct<(array<1 x ptr<1>>, array<1 x i64>)>, struct<(array<1 x ptr<1>>, array<1 x i64>)>, struct<(array<1 x ptr<1>>, array<1 x i64>)>, f32, f32, struct<(array<1 x ptr<1>>, array<1 x i64>)>, struct<(array<1 x ptr<1>>, array<1 x i64>)>, struct<(array<1 x ptr<1>>, array<1 x i64>)>, struct<(array<1 x ptr<1>>, array<1 x i64>)>, array<1 x ptr<1>>)>  loc(#loc)
    %16 = llvm.extractvalue %arg3[1, 2] : !llvm.struct<(array<1 x ptr<1>>, array<3 x i64>)>  loc(#loc)
    %17 = llvm.extractvalue %arg4[1, 2] : !llvm.struct<(array<1 x ptr<1>>, array<3 x i64>)>  loc(#loc)
    %18 = nvvm.read.ptx.sreg.tid.x : i32 loc(#loc16075)
    %19 = llvm.zext %18 : i32 to i64 loc(#loc17777)
    %20 = nvvm.read.ptx.sreg.ctaid.x : i32 loc(#loc16077)
    %21 = llvm.zext %20 : i32 to i64 loc(#loc17777)
    %22 = llvm.udiv %20, %0 : i32 loc(#loc19250)
    %23 = llvm.zext %22 : i32 to i64 loc(#loc19250)
    %24 = llvm.mul %23, %1 : i64 loc(#loc18426)

specifically

    %1 = llvm.mlir.constant(1152921504606846785 : i64) : i64 loc(#loc)
    %24 = llvm.mul %23, %1 : i64 loc(#loc18426)

this fancy mul, shouldn't have been produced [@gbaraldi this I believe was the reason for disabling instcombine]

[giordano]

Is it possible it's GPUCompiler? In Breeze.jl we're sticking to GPUCompiler@1.13.1 for the time being because something started causing us troubles in v1.13.2 with Reactant.

[gbaraldi]

I don't get what's fancy or weirdabout it. Is it that it messes with index math?

[wsmoses]

we've lost the actual logical indexing so conversion to affine indexing becomes busted. Aka someone did the trick to presumably turn a divide into a multiply by a fancy constant.

and yeah it's very possibly gpu compiler as the root cause for this second part

[gbaraldi]

I'm confused why not handle this in enzymejax? It doesn't look that weird. At least that much weirder than the original code probably would've been. Given we are dividing an index by 191

[wsmoses]

for the time being we assume the index math doesn't wrap around -- and that's what is breaking things here.

cuda.jl pre vendor IR
; ModuleID = 'start'
source_filename = "start"
target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
target triple = "nvptx64-nvidia-cuda"

%printf_args = type { i64, i32, i32, i32, i32, i32, i32 }
%printf_args.0 = type { i64 }
%printf_args.6 = type { i32, i64, i64, i32 }

@0 = private unnamed_addr constant [107 x i8] c"Stacktrace not available, run Julia on debug level 2 for more details (by passing -g2 to the executable).\0A\00", align 1
@1 = private unnamed_addr constant [4 x i8] c"%s\0A\00", align 1
@2 = private unnamed_addr constant [13 x i8] c"Stacktrace:\0A\00", align 1
@3 = private unnamed_addr constant [94 x i8] c"ERROR: a %s was thrown during kernel execution on thread (%d, %d, %d) in block (%d, %d, %d).\0A\00", align 1
@4 = private unnamed_addr constant [64 x i8] c"ERROR: Out of dynamic GPU memory (trying to allocate %d bytes)\0A\00", align 1
@5 = private unnamed_addr constant [2 x i8] c"\0A\00", align 1
@exception19 = private unnamed_addr constant [10 x i8] c"exception\00", align 1
@6 = private unnamed_addr constant [19 x i8] c" [%d] %s at %s:%d\0A\00", align 1

declare {}*** @julia.get_pgcstack() local_unnamed_addr

; Function Attrs: cold nounwind
declare void @llvm.trap() #0

; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #1

; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.y() #1

; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.z() #1

; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() #1

; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
declare i32 @llvm.nvvm.read.ptx.sreg.tid.y() #1

; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
declare i32 @llvm.nvvm.read.ptx.sreg.tid.z() #1

define ptx_kernel void @_Z29gpu__compute_barotropic_mode_16CompilerMetadataI10StaticSizeI8_48__24_E12DynamicCheckvv7NDRangeILi2ES0_I6_3__2_ES0_I8_16__16_EvvEE5FieldI4Face6Centervvvv11OffsetArrayI7Float32Li3E13CuTracedArrayISC_Li3ELi1E11_64__40__1_EESC_vvvES8_ISA_S9_vvvvSF_SC_vvvE21LatitudeLongitudeGridI15CuTracedRNumberISC_Li1EE8Periodic7BoundedSM_28StaticVerticalDiscretizationISB_ISC_Li1ESD_ISC_Li1ELi1E5_27__EESB_ISC_Li1ESD_ISC_Li1ELi1E5_26__EESP_SR_ESC_SC_SB_ISC_Li1E12StepRangeLenISC_7Float64SU_5Int64EESX_SC_SC_SX_SX_SB_ISC_Li1ESD_ISC_Li1ELi1E5_40__EESZ_SZ_SZ_SC_SC_vSV_ESB_ISC_Li3ESD_ISC_Li3ELi1E12_64__40__26_EES12_SF_({ { [1 x i8 addrspace(1)*], [3 x i64] } } %0, { { [1 x i8 addrspace(1)*], [3 x i64] } } %1, { i64, i64, i64, i64, i64, i64, [1 x i8 addrspace(1)*], [1 x i8 addrspace(1)*], [1 x i8 addrspace(1)*], float, float, { { double, double, i64, i64 }, [1 x i64] }, { { double, double, i64, i64 }, [1 x i64] }, float, float, { { double, double, i64, i64 }, [1 x i64] }, { { double, double, i64, i64 }, [1 x i64] }, [4 x { [1 x i8 addrspace(1)*], [1 x i64] }], { [1 x i8 addrspace(1)*], [1 x i64] }, { [1 x i8 addrspace(1)*], [1 x i64] }, { [1 x i8 addrspace(1)*], [1 x i64] }, { [1 x i8 addrspace(1)*], [1 x i64] }, float, float, { [1 x i8 addrspace(1)*], [1 x i64] }, { [1 x i8 addrspace(1)*], [1 x i64] }, { [1 x i8 addrspace(1)*], [1 x i64] }, { [1 x i8 addrspace(1)*], [1 x i64] }, [1 x i8 addrspace(1)*] } %2, { [1 x i8 addrspace(1)*], [3 x i64] } %3, { [1 x i8 addrspace(1)*], [3 x i64] } %4, { [1 x i8 addrspace(1)*], [3 x i64] } %5) local_unnamed_addr !dbg !39 {
conversion:
  %.fca.17.3.1.0.extract = extractvalue { i64, i64, i64, i64, i64, i64, [1 x i8 addrspace(1)*], [1 x i8 addrspace(1)*], [1 x i8 addrspace(1)*], float, float, { { double, double, i64, i64 }, [1 x i64] }, { { double, double, i64, i64 }, [1 x i64] }, float, float, { { double, double, i64, i64 }, [1 x i64] }, { { double, double, i64, i64 }, [1 x i64] }, [4 x { [1 x i8 addrspace(1)*], [1 x i64] }], { [1 x i8 addrspace(1)*], [1 x i64] }, { [1 x i8 addrspace(1)*], [1 x i64] }, { [1 x i8 addrspace(1)*], [1 x i64] }, { [1 x i8 addrspace(1)*], [1 x i64] }, float, float, { [1 x i8 addrspace(1)*], [1 x i64] }, { [1 x i8 addrspace(1)*], [1 x i64] }, { [1 x i8 addrspace(1)*], [1 x i64] }, { [1 x i8 addrspace(1)*], [1 x i64] }, [1 x i8 addrspace(1)*] } %2, 17, 3, 1, 0
  %.fca.1.2.extract15 = extractvalue { [1 x i8 addrspace(1)*], [3 x i64] } %3, 1, 2
  %.fca.1.2.extract5 = extractvalue { [1 x i8 addrspace(1)*], [3 x i64] } %4, 1, 2
  %6 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !dbg !43, !range !59
  %7 = zext i32 %6 to i64, !dbg !60
  %8 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x(), !dbg !78, !range !85
  %9 = zext i32 %8 to i64, !dbg !60
  %10 = udiv i32 %8, 3, !dbg !86
  %.zext = zext i32 %10 to i64, !dbg !86
  %.neg = mul i64 %.zext, 1152921504606846973, !dbg !92
  %11 = add i64 %.neg, %9, !dbg !95
  %12 = lshr i64 %7, 4, !dbg !86
  %13 = add nuw nsw i64 %12, 1, !dbg !96
  %14 = shl i64 %11, 4, !dbg !101
  %15 = shl nuw nsw i64 %.zext, 4, !dbg !101
  %16 = add nuw nsw i64 %13, %15, !dbg !108
  %17 = icmp ugt i64 %14, 47, !dbg !109
  %18 = icmp ugt i64 %16, 24, !dbg !122
  %.not50 = or i1 %18, %17, !dbg !58
  br i1 %.not50, label %L1200, label %pass6, !dbg !58

Note this is with Reactant.Compiler.DUMP_LLVMIR[] = true

and the first time it happens, so something is causing this within https://github.com/EnzymeAD/Reactant.jl/blob/321a9b3ed8481c36cffaeb5a1bc0552bb7f07586/ext/ReactantCUDAExt.jl#L913

[wsmoses]

function vendored_buildNewPMPipeline!(mpm, @nospecialize(job), opt_level)
    # Doesn't call instcombine
    GPUCompiler.buildEarlySimplificationPipeline(mpm, job, opt_level)
    LLVM.add!(mpm, LLVM.AlwaysInlinerPass())
    vendored_buildEarlyOptimizerPipeline(mpm, job, opt_level)
    LLVM.add!(mpm, LLVM.NewPMFunctionPassManager()) do fpm
        # Doesn't call instcombine
        GPUCompiler.buildLoopOptimizerPipeline(fpm, job, opt_level)
        vendored_buildScalarOptimizerPipeline(fpm, job, opt_level)
        if GPUCompiler.uses_julia_runtime(job) && opt_level >= 2
            # TODO(#2240): we disable vectorization, as this generally isn't useful for GPU targets
            #      and actually causes issues with some back-end compilers (like Metal).
            # TODO(#2240): Make this not dependent on `uses_julia_runtime` (likely CPU), but it's own control
            # Doesn't call instcombine
            GPUCompiler.buildVectorPipeline(fpm, job, opt_level)
        end     
        # if isdebug(:optim)
        #     add!(fpm, WarnMissedTransformationsPass())
        # end           
    end                 
    vendored_buildIntrinsicLoweringPipeline(mpm, job, opt_level)
    return GPUCompiler.buildCleanupPipeline(mpm, job, opt_level)
end  

yup, somebody added an instcombine =/ : https://github.com/JuliaGPU/GPUCompiler.jl/blame/8c79865568bcf7f79833b2393d910b2f3b3165ed/src/optim.jl#L135 in JuliaGPU/GPUCompiler.jl@c095147

[wsmoses]

@maleadt (cc @vchuravy) I didn't notice that your later changes actually removed this from an option within the target to a function on the job: (see initial pr version in JuliaGPU/GPUCompiler.jl@5b3f0ac#diff-62c39c7485e8353819ba216b11c1dd6c93879bbac92d57b122a6445a76085889)

I cannot fix this on our side without either:

• vendoring things more [defeating the point of the earlier pr]
• adding a parameter to the cudajob [which we can therefore specialize]
• type piracy on cuda's compiler job.

what's the right approach?