[Pytorch upstream] "div_rn" get "nan" result.

[etaf]

Describe the bug

Hi team:
Since PR pytorch/pytorch#164144 land, Inductor start generate triton.language.div_rn for float number div.
And I foud a series of Inductor UTs get nan result and failed, eg:

python test/inductor/test_compile_subprocess.py GPUTests.test_var_correction_xpu
python test/inductor/test_torchinductor.py GPUTests.test_unspec_inputs_int8_xpu
python test/inductor/test_torchinductor_dynamic_shapes.py DynamicShapesGPUTests.test_div_by_zero_dynamic_shapes_xpu
python test/inductor/test_compiled_optimizers.py CompiledOptimizerParityTestsXPU.test_correctness_RAdam_use_closure_True_xpu_float32

I've verifed that cuda works well.
You may start with test_div_by_zero_dynamic_shapes_xpu, which should get inf but nan.

Environment details

PyTorch: main branch (37d57ac9cb7f538b812cf1d9851b55b46213fe15)
Triton: pined
GPU: PVC

[anmyachev]

Simple Triton reproducer:

def test_div_rn(device):
    SIZE = 128

    @triton.jit
    def kernel(Z, X, Y, SIZE: tl.constexpr):
        off = tl.arange(0, SIZE)
        x = tl.load(X + off)
        y = tl.load(Y + off)
        z = tl.div_rn(x, y)
        tl.store(Z + off, z)

    x = torch.randn(10, device="xpu")
    y = torch.zeros(10, device="xpu")
    z = torch.zeros(10, device="xpu")
    kernel[(1, )](z, x, y, SIZE=SIZE, num_warps=4)
    print(z)  # <- tensor([nan, nan, nan, nan, nan, nan, nan, nan, nan, nan], device='xpu:0')
    print(x/y)  # <- tensor([-inf, -inf, inf, -inf, inf, -inf, -inf, -inf, -inf, inf], device='xpu:0')
    torch.testing.assert_close(z, x/y)

[anmyachev]

Blocked on intel/llvm#20427

[jinge90]

Blocked on intel/llvm#20427

Hi, @anmyachev
The fix has been merged, you can have a try.
Thanks very much.