Interpreter sort handles NaN values differently from XLA

[GleasonK]

What happened?

Note: This is a behavior difference, need to revisit IEEE to see if this classifies as a bug.

In updating testdata using XLA:CPU gold results, there seems to be a bug in sort (which is likely a bug in compare) when comparing NaN values.

module @jit_main attributes {mhlo.num_partitions = 1 : i32, mhlo.num_replicas = 1 : i32} {
  func.func public @main() -> (tensor<5xf32> {jax.result_info = "[0]", mhlo.layout_mode = "default"}, tensor<5xi32> {jax.result_info = "[1]", mhlo.layout_mode = "default"}) {
    %0 = call @inputs() : () -> tensor<5xf32>
    %1:2 = call @expected() : () -> (tensor<5xf32>, tensor<5xi32>)
    %2 = stablehlo.iota dim = 0 : tensor<5xi32>
    %3:2 = "stablehlo.sort"(%0, %2) <{dimension = 0 : i64, is_stable = true}> ({
    ^bb0(%arg0: tensor<f32>, %arg1: tensor<f32>, %arg2: tensor<i32>, %arg3: tensor<i32>):
      %6 = stablehlo.compare  GT, %arg0, %arg1,  TOTALORDER : (tensor<f32>, tensor<f32>) -> tensor<i1>
      stablehlo.return %6 : tensor<i1>
    }) : (tensor<5xf32>, tensor<5xi32>) -> (tensor<5xf32>, tensor<5xi32>)
    %4 = stablehlo.slice %3#0 [0:5] : (tensor<5xf32>) -> tensor<5xf32>
    %5 = stablehlo.slice %3#1 [0:5] : (tensor<5xi32>) -> tensor<5xi32>
    //stablehlo.custom_call @check.expect_eq(%4, %1#0) {has_side_effect = true} : (tensor<5xf32>, tensor<5xf32>) -> ()
    //stablehlo.custom_call @check.expect_eq(%5, %1#1) {has_side_effect = true} : (tensor<5xi32>, tensor<5xi32>) -> ()
    return %4, %5 : tensor<5xf32>, tensor<5xi32>
  }
  func.func private @inputs() -> (tensor<5xf32> {mhlo.layout_mode = "default"}) {
    //                              [+Inf,       +NaN,       -Inf,       -NaN,       3.0]
    %cst = stablehlo.constant dense<[0x7F800000, 0x7FC00000, 0xFFC00000, 0xFF800000, 3.000000e+00]> : tensor<5xf32>
    return %cst : tensor<5xf32>
  }
  func.func private @expected() -> (tensor<5xf32> {mhlo.layout_mode = "default"}, tensor<5xi32> {mhlo.layout_mode = "default"}) {
    // XLA:CPU
    //                              [+NaN,       +Inf,       3.0,          -Inf,       -NaN]
    %cst = stablehlo.constant dense<[0x7FC00000, 0x7F800000, 3.000000e+00, 0xFF800000, 0xFFC00000]> : tensor<5xf32>
    %c = stablehlo.constant dense<[1, 0, 4, 3, 2]> : tensor<5xi32>
    return %cst, %c : tensor<5xf32>, tensor<5xi32>
  }
  // StableHLO
  //   [+Inf,       +NaN,       -NaN,       3.0,          -Inf]
  //   [0x7F800000, 0x7FC00000, 0xFFC00000, 3.000000e+00, 0xFF800000]

This issue is being used to filter two interpreter tests, can search using RUN-DISABLED(#2497):.