Illegal frem error

cc @giordano @avik-pal

```mlir
(base) wmoses@hydra:~/git/Enzyme-JaX$ ./bazel-bin/enzymexlamlir-opt frem.mlir --lower-jit
:0: error: failed to legalize operation 'llvm.frem' that was explicitly marked illegal
:0: note: see current operation: %66 = "llvm.frem"(%63, %18) <{fastmathFlags = #llvm.fastmath}> : (f64, f64) -> f64
(base) wmoses@hydra:~/git/Enzyme-JaX$ cat frem.mlir
module attributes {gpu.container_module} {
func.func @main(%arg0: tensor<1x104x194xf64>, %arg1: tensor<1x104x194xf64>, %arg2: tensor<104x194xf64>, %arg3 : tensor<104x194xf64>) -> (tensor<1x104x194xf64>, tensor<1x104x194xf64>) {
%0:2 = enzymexla.jit_call @"##call__Z32gpu__compute_fractional_indices_16CompilerMetadataI16OffsetStaticSizeI13_0_181__0_91_E12DynamicCheckvv7NDRangeILi2E10StaticSizeI7_12__6_ES4_I8_16__16_E5TupleI5Int64S8_E13KernelOffsetsIS9_EEE10NamedTupleI8__i___j_S7_I5FieldI6CenterSG_vvvv11OffsetArrayI7Float64Li3E13CuTracedArrayISI_Li3ELi1E13_194__104__1_EESI_vvvESM_EE20ImmersedBoundaryGridISI_8Periodic14RightConnected7Bounded28OrthogonalSphericalShellGridISI_SQ_SR_SS_28StaticVerticalDiscretizationISH_ISI_Li1ESJ_ISI_Li1ELi1E5_35__EESH_ISI_Li1ESJ_ISI_Li1ELi1E5_34__EESW_SY_E8TripolarIS8_S8_S8_ESH_ISI_Li2ESJ_ISI_Li2ELi1E10_194__104_EES13_S13_S13_vE16GridFittedBottomISM_23CenterImmersedConditionEvvvE21LatitudeLongitudeGridISI_SQ_SS_4FlatSU_IvvSI_SI_ESI_SI_SH_ISI_Li1E12StepRangeLenISI_14TwicePrecisionISI_ES1E_S8_EES1G_SI_SI_S1G_S1G_SH_ISI_Li1ESJ_ISI_Li1ELi1E6_186__EES1I_S1I_S1I_SI_SI_vE#658$call$54" (%arg0, %arg1, %arg2, %arg3) {output_operand_aliases = [#stablehlo.output_operand_alias, #stablehlo.output_operand_alias]} : (tensor<1x104x194xf64>, tensor<1x104x194xf64>, tensor<104x194xf64>, tensor<104x194xf64>) -> (tensor<1x104x194xf64>, tensor<1x104x194xf64>)
return %0#0, %0#1 : tensor<1x104x194xf64>, tensor<1x104x194xf64>
}
gpu.module @gpumod___call__Z32gpu__compute_fractional_indices_16CompilerMetadataI16OffsetStaticSizeI13_0_181__0_91_E12DynamicCheckvv7NDRangeILi2E10StaticSizeI7_12__6_ES4_I8_16__16_E5TupleI5Int64S8_E13KernelOffsetsIS9_EEE10NamedTupleI8__i___j_S7_I5FieldI6CenterSG_vvvv11OffsetArrayI7Float64Li3E13CuTracedArrayISI_Li3ELi1E13_194__104__1_EESI_vvvESM_EE20ImmersedBoundaryGridISI_8Periodic14RightConnected7Bounded28OrthogonalSphericalShellGridISI_SQ_SR_SS_28StaticVerticalDiscretizationISH_ISI_Li1ESJ_ISI_Li1ELi1E5_35__EESH_ISI_Li1ESJ_ISI_Li1ELi1E5_34__EESW_SY_E8TripolarIS8_S8_S8_ESH_ISI_Li2ESJ_ISI_Li2ELi1E10_194__104_EES13_S13_S13_vE16GridFittedBottomISM_23CenterImmersedConditionEvvvE21LatitudeLongitudeGridISI_SQ_SS_4FlatSU_IvvSI_SI_ESI_SI_SH_ISI_Li1E12StepRangeLenISI_14TwicePrecisionISI_ES1E_S8_EES1G_SI_SI_S1G_S1G_SH_ISI_Li1ESJ_ISI_Li1ELi1E6_186__EES1I_S1I_S1I_SI_SI_vE_658 {
gpu.func @__call__Z32gpu__compute_fractional_indices_16CompilerMetadataI16OffsetStaticSizeI13_0_181__0_91_E12DynamicCheckvv7NDRangeILi2E10StaticSizeI7_12__6_ES4_I8_16__16_E5TupleI5Int64S8_E13KernelOffsetsIS9_EEE10NamedTupleI8__i___j_S7_I5FieldI6CenterSG_vvvv11OffsetArrayI7Float64Li3E13CuTracedArrayISI_Li3ELi1E13_194__104__1_EESI_vvvESM_EE20ImmersedBoundaryGridISI_8Periodic14RightConnected7Bounded28OrthogonalSphericalShellGridISI_SQ_SR_SS_28StaticVerticalDiscretizationISH_ISI_Li1ESJ_ISI_Li1ELi1E5_35__EESH_ISI_Li1ESJ_ISI_Li1ELi1E5_34__EESW_SY_E8TripolarIS8_S8_S8_ESH_ISI_Li2ESJ_ISI_Li2ELi1E10_194__104_EES13_S13_S13_vE16GridFittedBottomISM_23CenterImmersedConditionEvvvE21LatitudeLongitudeGridISI_SQ_SS_4FlatSU_IvvSI_SI_ESI_SI_SH_ISI_Li1E12StepRangeLenISI_14TwicePrecisionISI_ES1E_S8_EES1G_SI_SI_S1G_S1G_SH_ISI_Li1ESJ_ISI_Li1ELi1E6_186__EES1I_S1I_S1I_SI_SI_vE_658(%arg0: !llvm.ptr<1>, %arg1: !llvm.ptr<1>, %arg2: !llvm.ptr<1>, %arg3: !llvm.ptr<1>) kernel {
%false = arith.constant false
%cst = arith.constant -5.6843418860808015E-14 : f64
%cst_0 = arith.constant 5.000000e-01 : f64
%cst_1 = arith.constant -8.950000e+01 : f64
%c1_i32 = arith.constant 1 : i32
%c1_i64 = arith.constant 1 : i64
%c12_i32 = arith.constant 12 : i32
%c-12_i64 = arith.constant -12 : i64
%c16_i16 = arith.constant 16 : i16
%c0_i64 = arith.constant 0 : i64
%c16_i64 = arith.constant 16 : i64
%c181_i64 = arith.constant 181 : i64
%c91_i64 = arith.constant 91 : i64
%true = arith.constant true
%c6_i64 = arith.constant 6 : i64
%c194_i64 = arith.constant 194 : i64
%cst_2 = arith.constant 0.000000e+00 : f64
%cst_3 = arith.constant 1.000000e+00 : f64
%cst_4 = arith.constant 3.600000e+02 : f64
%c0_i8 = arith.constant 0 : i8
%0 = nvvm.read.ptx.sreg.ctaid.x range : i32
%1 = arith.addi %0, %c1_i32 : i32
%2 = nvvm.read.ptx.sreg.tid.x range : i32
%3 = arith.addi %2, %c1_i32 : i32
%4 = arith.extui %1 : i32 to i64
%5 = arith.subi %4, %c1_i64 : i64
%6 = arith.trunci %5 : i64 to i32
%7 = arith.divui %6, %c12_i32 : i32
%8 = arith.extui %7 : i32 to i64
%9 = arith.muli %8, %c-12_i64 : i64
%10 = arith.addi %5, %9 : i64
%11 = arith.extui %3 : i32 to i64
%12 = arith.subi %11, %c1_i64 : i64
%13 = arith.trunci %12 : i64 to i16
%14 = arith.divui %13, %c16_i16 : i16
%15 = arith.extui %14 : i16 to i64
%16 = arith.subi %c0_i64, %15 : i64
%17 = arith.addi %10, %16 : i64
%18 = arith.muli %17, %c16_i64 : i64
%19 = arith.addi %12, %18 : i64
%20 = arith.muli %8, %c16_i64 : i64
%21 = arith.addi %15, %20 : i64
%22 = arith.cmpi sge, %19, %c0_i64 : i64
%23 = arith.cmpi sle, %19, %c181_i64 : i64
%24 = arith.andi %22, %23 : i1
%25 = arith.cmpi sle, %21, %c91_i64 : i64
%26 = arith.andi %25, %24 : i1
%27 = arith.xori %26, %true : i1
llvm.cond_br %27, ^bb3, ^bb4
^bb1: // 2 preds: ^bb4, ^bb4
llvm.br ^bb2
^bb2: // pred: ^bb1
%28 = arith.cmpf ole, %cst_2, %55 {fastmathFlags = #llvm.fastmath} : f64
%29 = arith.select %28, %cst_2, %cst_3 : f64
%30 = arith.mulf %29, %cst_4 {fastmathFlags = #llvm.fastmath} : f64
%31 = arith.addf %55, %30 {fastmathFlags = #llvm.fastmath} : f64
%32 = arith.subf %31, %cst_0 {fastmathFlags = #llvm.fastmath} : f64
%33 = arith.addf %32, %cst_3 {fastmathFlags = #llvm.fastmath} : f64
%34 = arith.subf %45, %cst_1 {fastmathFlags = #llvm.fastmath} : f64
%35 = arith.addf %34, %cst_3 {fastmathFlags = #llvm.fastmath} : f64
%36 = llvm.getelementptr inbounds %arg0[%41] : (!llvm.ptr<1>, i64) -> !llvm.ptr<1>, f64
llvm.store %33, %36 {alignment = 8 : i64, tbaa = [#llvm.tbaa_tag, 0>}>, access_type = , 0>}>, offset = 0>]} : f64, !llvm.ptr<1>
%37 = llvm.getelementptr inbounds %arg1[%41] : (!llvm.ptr<1>, i64) -> !llvm.ptr<1>, f64
llvm.store %35, %37 {alignment = 8 : i64, tbaa = [#llvm.tbaa_tag, 0>}>, access_type = , 0>}>, offset = 0>]} : f64, !llvm.ptr<1>
llvm.br ^bb3
^bb3: // 2 preds: ^bb0, ^bb2
llvm.br ^bb5
^bb4: // pred: ^bb0
%38 = arith.addi %21, %c6_i64 : i64
%39 = arith.muli %38, %c194_i64 : i64
%40 = arith.addi %19, %c6_i64 : i64
%41 = arith.addi %40, %39 : i64
%42 = llvm.getelementptr inbounds %arg2[%41] : (!llvm.ptr<1>, i64) -> !llvm.ptr<1>, f64
%43 = llvm.load %42 {alignment = 8 : i64, tbaa = [#llvm.tbaa_tag, 0>}>, access_type = , 0>}>, offset = 0>]} : !llvm.ptr<1> -> f64
%44 = llvm.getelementptr inbounds %arg3[%41] : (!llvm.ptr<1>, i64) -> !llvm.ptr<1>, f64
%45 = llvm.load %44 {alignment = 8 : i64, tbaa = [#llvm.tbaa_tag, 0>}>, access_type = , 0>}>, offset = 0>]} : !llvm.ptr<1> -> f64
%46 = arith.remf %43, %cst_4 : f64
%47 = arith.addf %46, %cst_4 {fastmathFlags = #llvm.fastmath} : f64
%48 = arith.remf %47, %cst_4 : f64
%49 = arith.cmpf ole, %cst, %43 {fastmathFlags = #llvm.fastmath} : f64
%50 = arith.cmpf olt, %43, %cst_2 {fastmathFlags = #llvm.fastmath} : f64
%51 = arith.extui %50 : i1 to i8
%52 = arith.select %49, %51, %c0_i8 : i8
%53 = arith.addf %43, %cst_4 {fastmathFlags = #llvm.fastmath} : f64
%54 = arith.trunci %52 : i8 to i1
%55 = arith.select %54, %53, %48 : f64
llvm.cond_br %false, ^bb1, ^bb1
^bb5: // pred: ^bb3
gpu.return
}
}
func.func private @"##call__Z32gpu__compute_fractional_indices_16CompilerMetadataI16OffsetStaticSizeI13_0_181__0_91_E12DynamicCheckvv7NDRangeILi2E10StaticSizeI7_12__6_ES4_I8_16__16_E5TupleI5Int64S8_E13KernelOffsetsIS9_EEE10NamedTupleI8__i___j_S7_I5FieldI6CenterSG_vvvv11OffsetArrayI7Float64Li3E13CuTracedArrayISI_Li3ELi1E13_194__104__1_EESI_vvvESM_EE20ImmersedBoundaryGridISI_8Periodic14RightConnected7Bounded28OrthogonalSphericalShellGridISI_SQ_SR_SS_28StaticVerticalDiscretizationISH_ISI_Li1ESJ_ISI_Li1ELi1E5_35__EESH_ISI_Li1ESJ_ISI_Li1ELi1E5_34__EESW_SY_E8TripolarIS8_S8_S8_ESH_ISI_Li2ESJ_ISI_Li2ELi1E10_194__104_EES13_S13_S13_vE16GridFittedBottomISM_23CenterImmersedConditionEvvvE21LatitudeLongitudeGridISI_SQ_SS_4FlatSU_IvvSI_SI_ESI_SI_SH_ISI_Li1E12StepRangeLenISI_14TwicePrecisionISI_ES1E_S8_EES1G_SI_SI_S1G_S1G_SH_ISI_Li1ESJ_ISI_Li1ELi1E6_186__EES1I_S1I_S1I_SI_SI_vE#658$call$54"(%arg0: !llvm.ptr<1>, %arg1: !llvm.ptr<1>, %arg2: !llvm.ptr<1>, %arg3: !llvm.ptr<1>) {
%c72_i64 = arith.constant 72 : i64
%c1_i64 = arith.constant 1 : i64
%c256_i64 = arith.constant 256 : i64
%c0_i32 = arith.constant 0 : i32
%0 = "enzymexla.get_stream"() : () -> !gpu.async.token
%1 = gpu.launch_func async [%0] @gpumod___call__Z32gpu__compute_fractional_indices_16CompilerMetadataI16OffsetStaticSizeI13_0_181__0_91_E12DynamicCheckvv7NDRangeILi2E10StaticSizeI7_12__6_ES4_I8_16__16_E5TupleI5Int64S8_E13KernelOffsetsIS9_EEE10NamedTupleI8__i___j_S7_I5FieldI6CenterSG_vvvv11OffsetArrayI7Float64Li3E13CuTracedArrayISI_Li3ELi1E13_194__104__1_EESI_vvvESM_EE20ImmersedBoundaryGridISI_8Periodic14RightConnected7Bounded28OrthogonalSphericalShellGridISI_SQ_SR_SS_28StaticVerticalDiscretizationISH_ISI_Li1ESJ_ISI_Li1ELi1E5_35__EESH_ISI_Li1ESJ_ISI_Li1ELi1E5_34__EESW_SY_E8TripolarIS8_S8_S8_ESH_ISI_Li2ESJ_ISI_Li2ELi1E10_194__104_EES13_S13_S13_vE16GridFittedBottomISM_23CenterImmersedConditionEvvvE21LatitudeLongitudeGridISI_SQ_SS_4FlatSU_IvvSI_SI_ESI_SI_SH_ISI_Li1E12StepRangeLenISI_14TwicePrecisionISI_ES1E_S8_EES1G_SI_SI_S1G_S1G_SH_ISI_Li1ESJ_ISI_Li1ELi1E6_186__EES1I_S1I_S1I_SI_SI_vE_658::@__call__Z32gpu__compute_fractional_indices_16CompilerMetadataI16OffsetStaticSizeI13_0_181__0_91_E12DynamicCheckvv7NDRangeILi2E10StaticSizeI7_12__6_ES4_I8_16__16_E5TupleI5Int64S8_E13KernelOffsetsIS9_EEE10NamedTupleI8__i___j_S7_I5FieldI6CenterSG_vvvv11OffsetArrayI7Float64Li3E13CuTracedArrayISI_Li3ELi1E13_194__104__1_EESI_vvvESM_EE20ImmersedBoundaryGridISI_8Periodic14RightConnected7Bounded28OrthogonalSphericalShellGridISI_SQ_SR_SS_28StaticVerticalDiscretizationISH_ISI_Li1ESJ_ISI_Li1ELi1E5_35__EESH_ISI_Li1ESJ_ISI_Li1ELi1E5_34__EESW_SY_E8TripolarIS8_S8_S8_ESH_ISI_Li2ESJ_ISI_Li2ELi1E10_194__104_EES13_S13_S13_vE16GridFittedBottomISM_23CenterImmersedConditionEvvvE21LatitudeLongitudeGridISI_SQ_SS_4FlatSU_IvvSI_SI_ESI_SI_SH_ISI_Li1E12StepRangeLenISI_14TwicePrecisionISI_ES1E_S8_EES1G_SI_SI_S1G_S1G_SH_ISI_Li1ESJ_ISI_Li1ELi1E6_186__EES1I_S1I_S1I_SI_SI_vE_658 blocks in (%c72_i64, %c1_i64, %c1_i64) threads in (%c256_i64, %c1_i64, %c1_i64) : i64 dynamic_shared_memory_size %c0_i32 args(%arg0 : !llvm.ptr<1>, %arg1 : !llvm.ptr<1>, %arg2 : !llvm.ptr<1>, %arg3 : !llvm.ptr<1>)
return
}
}
```