[llvm] 309ae98 - [clang][NVVM][NVPTX] Remove pzo modifier for rs rounding mode conversions (#225012)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Sep 23 02:02:24 PDT 2026
Author: Srinivasa Ravi
Date: 2026-09-23T14:32:15+05:30
New Revision: 309ae98474b69d96541dc28dc6666117842a5f57
URL: https://github.com/llvm/llvm-project/commit/309ae98474b69d96541dc28dc6666117842a5f57
DIFF: https://github.com/llvm/llvm-project/commit/309ae98474b69d96541dc28dc6666117842a5f57.diff
LOG: [clang][NVVM][NVPTX] Remove pzo modifier for rs rounding mode conversions (#225012)
This change removes the `pzo` modifier from the `fp16x2` conversion
intrinsics with `rs` rounding mode. This was added in
https://github.com/llvm/llvm-project/pull/214667 due to a bug in the
developer-preview version of the 13.4 PTX ISA docs, and this change
aligns it with the latest PTX ISA docs for 13.4.
PTX ISA reference:
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cvt
Added:
Modified:
clang/test/CodeGen/builtins-nvptx.c
llvm/include/llvm/IR/IntrinsicsNVVM.td
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
llvm/test/CodeGen/NVPTX/convert-fp16-pzo.ll
llvm/test/CodeGen/NVPTX/convert-sm103a.ll
mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
mlir/test/Target/LLVMIR/nvvm/convert_fp16x2.mlir
Removed:
llvm/test/CodeGen/NVPTX/convert-fp16-rs-pzo.ll
################################################################################
diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c
index 607095bce481c..cd3eaff9d4c53 100644
--- a/clang/test/CodeGen/builtins-nvptx.c
+++ b/clang/test/CodeGen/builtins-nvptx.c
@@ -1394,51 +1394,51 @@ __device__ void nvvm_cvt_sm100a_sm103a() {
typedef __bf16 bf16x2 __attribute__((ext_vector_type(2)));
typedef char uint8x4 __attribute__((ext_vector_type(4)));
-// CHECK_PTX87_SM100a: %[[R1:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
+// CHECK_PTX87_SM100a: %[[R1:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0)
// CHECK_PTX87_SM100a: store <2 x half> %[[R1]], ptr %r1
-// CHECK_PTX87_SM103a: %[[R1:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
+// CHECK_PTX87_SM103a: %[[R1:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0)
// CHECK_PTX87_SM103a: store <2 x half> %[[R1]], ptr %r1
f16x2 r1 = __nvvm_ff2f16x2_rs(1.0f, 1.0f, 0);
-// CHECK_PTX87_SM100a: %[[R2:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
+// CHECK_PTX87_SM100a: %[[R2:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0)
// CHECK_PTX87_SM100a: store <2 x half> %[[R2]], ptr %r2
-// CHECK_PTX87_SM103a: %[[R2:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
+// CHECK_PTX87_SM103a: %[[R2:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0)
// CHECK_PTX87_SM103a: store <2 x half> %[[R2]], ptr %r2
f16x2 r2 = __nvvm_ff2f16x2_rs_relu(1.0f, 1.0f, 0);
-// CHECK_PTX87_SM100a: %[[R3:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
+// CHECK_PTX87_SM100a: %[[R3:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
// CHECK_PTX87_SM100a: store <2 x half> %[[R3]], ptr %r3
-// CHECK_PTX87_SM103a: %[[R3:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
+// CHECK_PTX87_SM103a: %[[R3:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
// CHECK_PTX87_SM103a: store <2 x half> %[[R3]], ptr %r3
f16x2 r3 = __nvvm_ff2f16x2_rs_satfinite(1.0f, 1.0f, 0);
-// CHECK_PTX87_SM100a: %[[R4:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
+// CHECK_PTX87_SM100a: %[[R4:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
// CHECK_PTX87_SM100a: store <2 x half> %[[R4]], ptr %r4
-// CHECK_PTX87_SM103a: %[[R4:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
+// CHECK_PTX87_SM103a: %[[R4:.*]] = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
// CHECK_PTX87_SM103a: store <2 x half> %[[R4]], ptr %r4
f16x2 r4 = __nvvm_ff2f16x2_rs_relu_satfinite(1.0f, 1.0f, 0);
-// CHECK_PTX87_SM100a: %[[R5:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
+// CHECK_PTX87_SM100a: %[[R5:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0)
// CHECK_PTX87_SM100a: store <2 x bfloat> %[[R5]], ptr %r5
-// CHECK_PTX87_SM103a: %[[R5:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
+// CHECK_PTX87_SM103a: %[[R5:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float 1.000000e+00, float 1.000000e+00, i32 0)
// CHECK_PTX87_SM103a: store <2 x bfloat> %[[R5]], ptr %r5
bf16x2 r5 = __nvvm_ff2bf16x2_rs(1.0f, 1.0f, 0);
-// CHECK_PTX87_SM100a: %[[R6:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
+// CHECK_PTX87_SM100a: %[[R6:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0)
// CHECK_PTX87_SM100a: store <2 x bfloat> %[[R6]], ptr %r6
-// CHECK_PTX87_SM103a: %[[R6:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
+// CHECK_PTX87_SM103a: %[[R6:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float 1.000000e+00, float 1.000000e+00, i32 0)
// CHECK_PTX87_SM103a: store <2 x bfloat> %[[R6]], ptr %r6
bf16x2 r6 = __nvvm_ff2bf16x2_rs_relu(1.0f, 1.0f, 0);
-// CHECK_PTX87_SM100a: %[[R7:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
+// CHECK_PTX87_SM100a: %[[R7:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
// CHECK_PTX87_SM100a: store <2 x bfloat> %[[R7]], ptr %r7
-// CHECK_PTX87_SM103a: %[[R7:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
+// CHECK_PTX87_SM103a: %[[R7:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
// CHECK_PTX87_SM103a: store <2 x bfloat> %[[R7]], ptr %r7
bf16x2 r7 = __nvvm_ff2bf16x2_rs_satfinite(1.0f, 1.0f, 0);
-// CHECK_PTX87_SM100a: %[[R8:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
+// CHECK_PTX87_SM100a: %[[R8:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
// CHECK_PTX87_SM100a: store <2 x bfloat> %[[R8]], ptr %r8
-// CHECK_PTX87_SM103a: %[[R8:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0, i1 false)
+// CHECK_PTX87_SM103a: %[[R8:.*]] = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float 1.000000e+00, float 1.000000e+00, i32 0)
// CHECK_PTX87_SM103a: store <2 x bfloat> %[[R8]], ptr %r8
bf16x2 r8 = __nvvm_ff2bf16x2_rs_relu_satfinite(1.0f, 1.0f, 0);
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index d9b73418394ed..95e1def383428 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2065,10 +2065,10 @@ let TargetPrefix = "nvvm" in {
foreach relu = ["", "_relu"] in {
foreach satfinite = ["", "_satfinite"] in {
def int_nvvm_ff2f16x2_rs # relu # satfinite : NVVMBuiltin,
- PureIntrinsic<[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i1_ty], [ImmArg<ArgIndex<3>, DefaultValue<0>>]>;
+ PureIntrinsic<[llvm_v2f16_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty]>;
def int_nvvm_ff2bf16x2_rs # relu # satfinite : NVVMBuiltin,
- PureIntrinsic<[llvm_v2bf16_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty, llvm_i1_ty], [ImmArg<ArgIndex<3>, DefaultValue<0>>]>;
+ PureIntrinsic<[llvm_v2bf16_ty], [llvm_float_ty, llvm_float_ty, llvm_i32_ty]>;
}
}
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index f794002971905..3fa66e774ae61 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -741,13 +741,13 @@ let hasSideEffects = false in {
BasicFlagsNVPTXInst<(outs RC:$dst),
(ins B32:$src1, B32:$src2, B32:$src3),
(ins CvtMode:$mode),
- "cvt${mode:base}${mode:relu}${mode:pzo}." # FromName # ".f32">;
+ "cvt${mode:base}${mode:relu}." # FromName # ".f32">;
def _f32_rs_sf :
BasicFlagsNVPTXInst<(outs RC:$dst),
(ins B32:$src1, B32:$src2, B32:$src3),
(ins CvtMode:$mode),
- "cvt${mode:base}${mode:relu}.satfinite${mode:pzo}." # FromName # ".f32">;
+ "cvt${mode:base}${mode:relu}.satfinite." # FromName # ".f32">;
}
defm CVT_f16x2 : CVT_FROM_FLOAT_V2_RS<"f16x2", B32>;
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 89a9adf4426d8..7b0848cef1186 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -3027,28 +3027,21 @@ foreach rnd = ["rn", "rz"] in {
}
}
-foreach relu = ["", "_relu"] in {
- defvar BaseMode = !cast<PatLeaf>("CvtRS" # !toupper(relu));
-
- foreach pzo = [0, 1] in {
- defvar PZO = !if(pzo, -1, 0);
- defvar Mode = CvtModeWithPZO<BaseMode, pzo>.Value;
- defvar Preds =
- !listconcat(
- [hasConvertWithStochasticRounding]<Predicate>,
- !if(pzo, [hasConvertWithPZOSupport]<Predicate>, []<Predicate>));
+let Predicates = [hasConvertWithStochasticRounding] in {
+ foreach relu = ["", "_relu"] in {
+ defvar Mode = !cast<PatLeaf>("CvtRS" # !toupper(relu));
defvar BF16X2 = ToCvtIntrinsics<"ff2bf16x2_rs" # relu>;
defvar F16X2 = ToCvtIntrinsics<"ff2f16x2_rs" # relu>;
- def : Pat<(BF16X2.base f32:$a, f32:$b, i32:$c, PZO),
- (CVT_bf16x2_f32_rs $a, $b, $c, Mode)>, Requires<Preds>;
- def : Pat<(BF16X2.sf f32:$a, f32:$b, i32:$c, PZO),
- (CVT_bf16x2_f32_rs_sf $a, $b, $c, Mode)>, Requires<Preds>;
- def : Pat<(F16X2.base f32:$a, f32:$b, i32:$c, PZO),
- (CVT_f16x2_f32_rs $a, $b, $c, Mode)>, Requires<Preds>;
- def : Pat<(F16X2.sf f32:$a, f32:$b, i32:$c, PZO),
- (CVT_f16x2_f32_rs_sf $a, $b, $c, Mode)>, Requires<Preds>;
+ def : Pat<(BF16X2.base f32:$a, f32:$b, i32:$c),
+ (CVT_bf16x2_f32_rs $a, $b, $c, Mode)>;
+ def : Pat<(BF16X2.sf f32:$a, f32:$b, i32:$c),
+ (CVT_bf16x2_f32_rs_sf $a, $b, $c, Mode)>;
+ def : Pat<(F16X2.base f32:$a, f32:$b, i32:$c),
+ (CVT_f16x2_f32_rs $a, $b, $c, Mode)>;
+ def : Pat<(F16X2.sf f32:$a, f32:$b, i32:$c),
+ (CVT_f16x2_f32_rs_sf $a, $b, $c, Mode)>;
}
}
diff --git a/llvm/test/CodeGen/NVPTX/convert-fp16-pzo.ll b/llvm/test/CodeGen/NVPTX/convert-fp16-pzo.ll
index 6861daf25aacf..339984387d9ca 100644
--- a/llvm/test/CodeGen/NVPTX/convert-fp16-pzo.ll
+++ b/llvm/test/CodeGen/NVPTX/convert-fp16-pzo.ll
@@ -2,8 +2,8 @@
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_107f -mattr=+ptx94 | FileCheck %s
; RUN: %if ptxas-sm_107f && ptxas-isa-9.4 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_107f -mattr=+ptx94 | %ptxas-verify -arch=sm_107f %}
-; Conversions to fp16 types with the preserve-zero-order modifier. The RS
-; (stochastic rounding) variants need sm_107a and live in convert-fp16-rs-pzo.ll.
+; Conversions to fp16 types with the preserve-zero-order modifier. The modifier
+; is not supported with the RS (stochastic rounding) rounding mode.
; F16X2 conversions
diff --git a/llvm/test/CodeGen/NVPTX/convert-fp16-rs-pzo.ll b/llvm/test/CodeGen/NVPTX/convert-fp16-rs-pzo.ll
deleted file mode 100644
index 58eacc36bcf55..0000000000000
--- a/llvm/test/CodeGen/NVPTX/convert-fp16-rs-pzo.ll
+++ /dev/null
@@ -1,135 +0,0 @@
-; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
-; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_107a -mattr=+ptx94 | FileCheck %s
-; RUN: %if ptxas-sm_107a && ptxas-isa-9.4 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_107a -mattr=+ptx94 | %ptxas-verify -arch=sm_107a %}
-
-; F16X2 conversions
-
-define <2 x half> @cvt_rs_pzo_f16x2_f32(float %f1, float %f2, i32 %rbits) {
-; CHECK-LABEL: cvt_rs_pzo_f16x2_f32(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<5>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_pzo_f16x2_f32_param_0];
-; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_pzo_f16x2_f32_param_1];
-; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_pzo_f16x2_f32_param_2];
-; CHECK-NEXT: cvt.rs.pzo.f16x2.f32 %r4, %r1, %r2, %r3;
-; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4;
-; CHECK-NEXT: ret;
- %val = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float %f1, float %f2, i32 %rbits, i1 true)
- ret <2 x half> %val
-}
-
-define <2 x half> @cvt_rs_relu_pzo_f16x2_f32(float %f1, float %f2, i32 %rbits) {
-; CHECK-LABEL: cvt_rs_relu_pzo_f16x2_f32(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<5>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_relu_pzo_f16x2_f32_param_0];
-; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_relu_pzo_f16x2_f32_param_1];
-; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_relu_pzo_f16x2_f32_param_2];
-; CHECK-NEXT: cvt.rs.relu.pzo.f16x2.f32 %r4, %r1, %r2, %r3;
-; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4;
-; CHECK-NEXT: ret;
- %val = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float %f1, float %f2, i32 %rbits, i1 true)
- ret <2 x half> %val
-}
-
-define <2 x half> @cvt_rs_sf_pzo_f16x2_f32(float %f1, float %f2, i32 %rbits) {
-; CHECK-LABEL: cvt_rs_sf_pzo_f16x2_f32(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<5>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_sf_pzo_f16x2_f32_param_0];
-; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_sf_pzo_f16x2_f32_param_1];
-; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_sf_pzo_f16x2_f32_param_2];
-; CHECK-NEXT: cvt.rs.satfinite.pzo.f16x2.f32 %r4, %r1, %r2, %r3;
-; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4;
-; CHECK-NEXT: ret;
- %val = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float %f1, float %f2, i32 %rbits, i1 true)
- ret <2 x half> %val
-}
-
-define <2 x half> @cvt_rs_relu_sf_pzo_f16x2_f32(float %f1, float %f2, i32 %rbits) {
-; CHECK-LABEL: cvt_rs_relu_sf_pzo_f16x2_f32(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<5>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_relu_sf_pzo_f16x2_f32_param_0];
-; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_relu_sf_pzo_f16x2_f32_param_1];
-; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_relu_sf_pzo_f16x2_f32_param_2];
-; CHECK-NEXT: cvt.rs.relu.satfinite.pzo.f16x2.f32 %r4, %r1, %r2, %r3;
-; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4;
-; CHECK-NEXT: ret;
- %val = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float %f1, float %f2, i32 %rbits, i1 true)
- ret <2 x half> %val
-}
-
-; BF16X2 conversions
-
-define <2 x bfloat> @cvt_rs_pzo_bf16x2_f32(float %f1, float %f2, i32 %rbits) {
-; CHECK-LABEL: cvt_rs_pzo_bf16x2_f32(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<5>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_pzo_bf16x2_f32_param_0];
-; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_pzo_bf16x2_f32_param_1];
-; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_pzo_bf16x2_f32_param_2];
-; CHECK-NEXT: cvt.rs.pzo.bf16x2.f32 %r4, %r1, %r2, %r3;
-; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4;
-; CHECK-NEXT: ret;
- %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float %f1, float %f2, i32 %rbits, i1 true)
- ret <2 x bfloat> %val
-}
-
-define <2 x bfloat> @cvt_rs_relu_pzo_bf16x2_f32(float %f1, float %f2, i32 %rbits) {
-; CHECK-LABEL: cvt_rs_relu_pzo_bf16x2_f32(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<5>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_relu_pzo_bf16x2_f32_param_0];
-; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_relu_pzo_bf16x2_f32_param_1];
-; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_relu_pzo_bf16x2_f32_param_2];
-; CHECK-NEXT: cvt.rs.relu.pzo.bf16x2.f32 %r4, %r1, %r2, %r3;
-; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4;
-; CHECK-NEXT: ret;
- %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float %f1, float %f2, i32 %rbits, i1 true)
- ret <2 x bfloat> %val
-}
-
-define <2 x bfloat> @cvt_rs_sf_pzo_bf16x2_f32(float %f1, float %f2, i32 %rbits) {
-; CHECK-LABEL: cvt_rs_sf_pzo_bf16x2_f32(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<5>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_sf_pzo_bf16x2_f32_param_0];
-; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_sf_pzo_bf16x2_f32_param_1];
-; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_sf_pzo_bf16x2_f32_param_2];
-; CHECK-NEXT: cvt.rs.satfinite.pzo.bf16x2.f32 %r4, %r1, %r2, %r3;
-; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4;
-; CHECK-NEXT: ret;
- %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float %f1, float %f2, i32 %rbits, i1 true)
- ret <2 x bfloat> %val
-}
-
-define <2 x bfloat> @cvt_rs_relu_sf_pzo_bf16x2_f32(float %f1, float %f2, i32 %rbits) {
-; CHECK-LABEL: cvt_rs_relu_sf_pzo_bf16x2_f32(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<5>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param::func.b32 %r1, [cvt_rs_relu_sf_pzo_bf16x2_f32_param_0];
-; CHECK-NEXT: ld.param::func.b32 %r2, [cvt_rs_relu_sf_pzo_bf16x2_f32_param_1];
-; CHECK-NEXT: ld.param::func.b32 %r3, [cvt_rs_relu_sf_pzo_bf16x2_f32_param_2];
-; CHECK-NEXT: cvt.rs.relu.satfinite.pzo.bf16x2.f32 %r4, %r1, %r2, %r3;
-; CHECK-NEXT: st.param::func.b32 [func_retval0], %r4;
-; CHECK-NEXT: ret;
- %val = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float %f1, float %f2, i32 %rbits, i1 true)
- ret <2 x bfloat> %val
-}
diff --git a/llvm/test/CodeGen/NVPTX/convert-sm103a.ll b/llvm/test/CodeGen/NVPTX/convert-sm103a.ll
index 75a394e4764ed..6a452d1d66baa 100644
--- a/llvm/test/CodeGen/NVPTX/convert-sm103a.ll
+++ b/llvm/test/CodeGen/NVPTX/convert-sm103a.ll
@@ -1,8 +1,10 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100a -mattr=+ptx87 | FileCheck %s
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | FileCheck %s
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_107a -mattr=+ptx94 | FileCheck %s
; RUN: %if ptxas-sm_100a && ptxas-isa-8.7 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100a -mattr=+ptx87 | %ptxas-verify -arch=sm_100a %}
; RUN: %if ptxas-sm_103a && ptxas-isa-8.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_103a -mattr=+ptx88 | %ptxas-verify -arch=sm_103a %}
+; RUN: %if ptxas-sm_107a && ptxas-isa-9.4 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_107a -mattr=+ptx94 | %ptxas-verify -arch=sm_107a %}
; F16X2 conversions
diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
index 46e9ab29bc4a0..f9f8a3c4387d7 100644
--- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
+++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp
@@ -5674,8 +5674,10 @@ ConvertF32x2ToF16x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToF16x2Op &op,
if (op.getRandomBits())
args.push_back(mt.lookupValue(op.getRandomBits()));
+ // The PZO modifier is not supported with the RS rounding mode.
// TODO: Add support for PZO modifier
- args.push_back(builder.getInt1(false));
+ if (op.getRnd() != FPRoundingMode::RS)
+ args.push_back(builder.getInt1(false));
switch (op.getRnd()) {
case FPRoundingMode::RN:
@@ -5726,7 +5728,8 @@ ConvertF32x2ToBF16x2Op::getIntrinsicIDAndArgs(NVVM::ConvertF32x2ToBF16x2Op &op,
args.push_back(mt.lookupValue(op.getRandomBits()));
// TODO: Add support for PZO modifier
- args.push_back(builder.getInt1(false));
+ if (op.getRnd() != FPRoundingMode::RS)
+ args.push_back(builder.getInt1(false));
switch (op.getRnd()) {
case FPRoundingMode::RN:
diff --git a/mlir/test/Target/LLVMIR/nvvm/convert_fp16x2.mlir b/mlir/test/Target/LLVMIR/nvvm/convert_fp16x2.mlir
index ddb6defe270bb..6338f6a15030e 100644
--- a/mlir/test/Target/LLVMIR/nvvm/convert_fp16x2.mlir
+++ b/mlir/test/Target/LLVMIR/nvvm/convert_fp16x2.mlir
@@ -30,13 +30,13 @@ llvm.func @convert_f32x2_to_f16x2_rz(%srcA : f32, %srcB : f32) {
// CHECK-LABEL: @convert_f32x2_to_f16x2_rs_stochastic
llvm.func @convert_f32x2_to_f16x2_rs_stochastic(%srcA : f32, %srcB : f32, %rbits : i32) {
- // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false)
+ // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs(float %{{.*}}, float %{{.*}}, i32 %{{.*}})
%res1 = nvvm.convert.f32x2.to.f16x2 %srcA, %srcB random_bits = %rbits rnd = <rs> : vector<2xf16>
- // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false)
+ // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu(float %{{.*}}, float %{{.*}}, i32 %{{.*}})
%res2 = nvvm.convert.f32x2.to.f16x2 %srcA, %srcB random_bits = %rbits rnd = <rs> relu = true : vector<2xf16>
- // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false)
+ // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs.satfinite(float %{{.*}}, float %{{.*}}, i32 %{{.*}})
%res3 = nvvm.convert.f32x2.to.f16x2 %srcA, %srcB random_bits = %rbits rnd = <rs> sat = <satfinite> : vector<2xf16>
- // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false)
+ // CHECK: %{{.*}} = call <2 x half> @llvm.nvvm.ff2f16x2.rs.relu.satfinite(float %{{.*}}, float %{{.*}}, i32 %{{.*}})
%res4 = nvvm.convert.f32x2.to.f16x2 %srcA, %srcB random_bits = %rbits rnd = <rs> sat = <satfinite> relu = true : vector<2xf16>
llvm.return
@@ -74,13 +74,13 @@ llvm.func @convert_f32x2_to_bf16x2_rz(%srcA : f32, %srcB : f32) {
// CHECK-LABEL: @convert_f32x2_to_bf16x2_rs_stochastic
llvm.func @convert_f32x2_to_bf16x2_rs_stochastic(%srcA : f32, %srcB : f32, %rbits : i32) {
- // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false)
+ // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs(float %{{.*}}, float %{{.*}}, i32 %{{.*}})
%res1 = nvvm.convert.f32x2.to.bf16x2 %srcA, %srcB random_bits = %rbits rnd = <rs> : vector<2xbf16>
- // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false)
+ // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu(float %{{.*}}, float %{{.*}}, i32 %{{.*}})
%res2 = nvvm.convert.f32x2.to.bf16x2 %srcA, %srcB random_bits = %rbits rnd = <rs> relu = true : vector<2xbf16>
- // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false)
+ // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.satfinite(float %{{.*}}, float %{{.*}}, i32 %{{.*}})
%res3 = nvvm.convert.f32x2.to.bf16x2 %srcA, %srcB random_bits = %rbits rnd = <rs> sat = <satfinite> : vector<2xbf16>
- // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float %{{.*}}, float %{{.*}}, i32 %{{.*}}, i1 false)
+ // CHECK: %{{.*}} = call <2 x bfloat> @llvm.nvvm.ff2bf16x2.rs.relu.satfinite(float %{{.*}}, float %{{.*}}, i32 %{{.*}})
%res4 = nvvm.convert.f32x2.to.bf16x2 %srcA, %srcB random_bits = %rbits rnd = <rs> sat = <satfinite> relu = true : vector<2xbf16>
llvm.return
More information about the llvm-commits
mailing list