[NVPTX] Add support for integer min/max ReLU idiom #151727

AlexMaclean · 2025-08-01T16:40:57Z

No description provided.

llvmbot · 2025-08-01T16:41:41Z

@llvm/pr-subscribers-backend-nvptx

Author: Alex MacLean (AlexMaclean)

Changes

Patch is 35.24 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/151727.diff

2 Files Affected:

(modified) llvm/lib/Target/NVPTX/NVPTXInstrInfo.td (+21-5)
(modified) llvm/test/CodeGen/NVPTX/combine-min-max.ll (+629-104)

diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index d8047d31ff6f0..e42fca6b3b2a1 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -151,6 +151,8 @@ class OneUse2<SDPatternOperator operator>
 class fpimm_pos_inf<ValueType vt>
     : FPImmLeaf<vt, [{ return Imm.isPosInfinity(); }]>;
 
+class zeroinitializer<ValueType vt> : 
+  PatLeaf<(vt (bitconvert (!cast<ValueType>("i" # vt.Size) 0)))>;
 
 
 // Operands which can hold a Register or an Immediate.
@@ -789,6 +791,23 @@ def UMAX16x2 : I16x2<"max.u", umax>;
 def SMIN16x2 : I16x2<"min.s", smin>;
 def UMIN16x2 : I16x2<"min.u", umin>;
 
+let Predicates = [hasPTX<80>, hasSM<90>, hasOptEnabled] in {
+
+  def MIN_RELU_S32 : BasicNVPTXInst<(outs B32:$dst), (ins B32:$a, B32:$b),
+                     "min.relu.s32",
+                     [(set i32:$dst, (smax (smin i32:$a, i32:$b), 0))]>;
+  def MAX_RELU_S32 : BasicNVPTXInst<(outs B32:$dst), (ins B32:$a, B32:$b),
+                     "max.relu.s32",
+                     [(set i32:$dst, (smax (smax i32:$a, i32:$b), 0))]>;
+  def MIN_RELU_S16x2 : BasicNVPTXInst<(outs B32:$dst), (ins B32:$a, B32:$b),
+                     "min.relu.s16x2",
+                     [(set v2i16:$dst, (smax (smin v2i16:$a, v2i16:$b),
+                                             zeroinitializer<v2i16>))]>;
+  def MAX_RELU_S16x2 : BasicNVPTXInst<(outs B32:$dst), (ins B32:$a, B32:$b),
+                     "max.relu.s16x2",
+                     [(set v2i16:$dst, (smax (smax v2i16:$a, v2i16:$b),
+                                             zeroinitializer<v2i16>))]>;
+}
 
 //
 // Wide multiplication
@@ -2385,9 +2404,6 @@ def fpimm_any_zero : FPImmLeaf<fAny, [{
   return Imm.isZero();
 }]>;
 
-def fpimm_positive_zero_v2f16 : PatFrag<(ops), (v2f16 (bitconvert (i32 0)))>;
-def fpimm_positive_zero_v2bf16 : PatFrag<(ops), (v2bf16 (bitconvert (i32 0)))>;
-
 // Perform substitution if fma only has one use, and also if instruction has
 // nnan instruction flag or if the TM has NoNaNsFPMath
 def NVPTX_fma_oneuse_and_nnan : PatFrag<(ops node:$a, node:$b, node:$c),
@@ -2410,10 +2426,10 @@ class FMARELUInst<RegTyInfo t, bit allow_ftz, PatFrag zero_pat>
 
 let Predicates = [useFP16Math, hasPTX<70>, hasSM<80>] in {
   def FMARELU_F16 : FMARELUInst<F16RT, true, fpimm_any_zero>;
-  def FMARELU_F16X2 : FMARELUInst<F16X2RT, true, fpimm_positive_zero_v2f16>;
+  def FMARELU_F16X2 : FMARELUInst<F16X2RT, true, zeroinitializer<v2f16>>;
 }
 
 let Predicates = [hasBF16Math, hasPTX<70>, hasSM<80>] in {
   def FMARELU_BF16 : FMARELUInst<BF16RT, false, fpimm_any_zero>;
-  def FMARELU_BF16X2 : FMARELUInst<BF16X2RT, false, fpimm_positive_zero_v2bf16>;
+  def FMARELU_BF16X2 : FMARELUInst<BF16X2RT, false, zeroinitializer<v2bf16>>;
 }
diff --git a/llvm/test/CodeGen/NVPTX/combine-min-max.ll b/llvm/test/CodeGen/NVPTX/combine-min-max.ll
index 157c3cc6bd2e4..a18aa77a705b9 100644
--- a/llvm/test/CodeGen/NVPTX/combine-min-max.ll
+++ b/llvm/test/CodeGen/NVPTX/combine-min-max.ll
@@ -1,22 +1,35 @@
-; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O2 | FileCheck %s
-; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 -O2 | %ptxas-verify %}
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -O3 | FileCheck %s
+; RUN: %if ptxas-12.0 %{ llc < %s -mcpu=sm_90 -mattr=+ptx80 -O3 | %ptxas-verify %}
+
+target triple = "nvptx64-nvidia-cuda"
 
 ; *************************************
 ; * Cases with no min/max
 
 define i32 @ab_eq_i32(i32 %a, i32 %b) {
-; CHECK-LABEL: @ab_eq_i32
-; CHECK-NOT: min
-; CHECK-NOT: max
+; CHECK-LABEL: ab_eq_i32(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b32 %r1, [ab_eq_i32_param_1];
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
   %cmp = icmp eq i32 %a, %b
   %sel = select i1 %cmp, i32 %a, i32 %b
   ret i32 %sel
 }
 
 define i64 @ab_ne_i64(i64 %a, i64 %b) {
-; CHECK-LABEL: @ab_ne_i64
-; CHECK-NOT: min
-; CHECK-NOT: max
+; CHECK-LABEL: ab_ne_i64(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b64 %rd1, [ab_ne_i64_param_1];
+; CHECK-NEXT:    st.param.b64 [func_retval0], %rd1;
+; CHECK-NEXT:    ret;
   %cmp = icmp ne i64 %a, %b
   %sel = select i1 %cmp, i64 %b, i64 %a
   ret i64 %sel
@@ -27,32 +40,72 @@ define i64 @ab_ne_i64(i64 %a, i64 %b) {
 
 ; *** ab, unsigned, i16
 define i16 @ab_ugt_i16(i16 %a, i16 %b) {
-; CHECK-LABEL: @ab_ugt_i16
-; CHECK: max.u16
+; CHECK-LABEL: ab_ugt_i16(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<4>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b16 %rs1, [ab_ugt_i16_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs2, [ab_ugt_i16_param_1];
+; CHECK-NEXT:    max.u16 %rs3, %rs1, %rs2;
+; CHECK-NEXT:    cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
   %cmp = icmp ugt i16 %a, %b
   %sel = select i1 %cmp, i16 %a, i16 %b
   ret i16 %sel
 }
 
 define i16 @ab_uge_i16(i16 %a, i16 %b) {
-; CHECK-LABEL: @ab_uge_i16
-; CHECK: max.u16
+; CHECK-LABEL: ab_uge_i16(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<4>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b16 %rs1, [ab_uge_i16_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs2, [ab_uge_i16_param_1];
+; CHECK-NEXT:    max.u16 %rs3, %rs1, %rs2;
+; CHECK-NEXT:    cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
   %cmp = icmp uge i16 %a, %b
   %sel = select i1 %cmp, i16 %a, i16 %b
   ret i16 %sel
 }
 
 define i16 @ab_ult_i16(i16 %a, i16 %b) {
-; CHECK-LABEL: @ab_ult_i16
-; CHECK: min.u16
+; CHECK-LABEL: ab_ult_i16(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<4>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b16 %rs1, [ab_ult_i16_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs2, [ab_ult_i16_param_1];
+; CHECK-NEXT:    min.u16 %rs3, %rs1, %rs2;
+; CHECK-NEXT:    cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
   %cmp = icmp ult i16 %a, %b
   %sel = select i1 %cmp, i16 %a, i16 %b
   ret i16 %sel
 }
 
 define i16 @ab_ule_i16(i16 %a, i16 %b) {
-; CHECK-LABEL: @ab_ule_i16
-; CHECK: min.u16
+; CHECK-LABEL: ab_ule_i16(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<4>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b16 %rs1, [ab_ule_i16_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs2, [ab_ule_i16_param_1];
+; CHECK-NEXT:    min.u16 %rs3, %rs1, %rs2;
+; CHECK-NEXT:    cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
   %cmp = icmp ule i16 %a, %b
   %sel = select i1 %cmp, i16 %a, i16 %b
   ret i16 %sel
@@ -60,32 +113,72 @@ define i16 @ab_ule_i16(i16 %a, i16 %b) {
 
 ; *** ab, signed, i16
 define i16 @ab_sgt_i16(i16 %a, i16 %b) {
-; CHECK-LABEL: @ab_sgt_i16
-; CHECK: max.s16
+; CHECK-LABEL: ab_sgt_i16(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<4>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b16 %rs1, [ab_sgt_i16_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs2, [ab_sgt_i16_param_1];
+; CHECK-NEXT:    max.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT:    cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
   %cmp = icmp sgt i16 %a, %b
   %sel = select i1 %cmp, i16 %a, i16 %b
   ret i16 %sel
 }
 
 define i16 @ab_sge_i16(i16 %a, i16 %b) {
-; CHECK-LABEL: @ab_sge_i16
-; CHECK: max.s16
+; CHECK-LABEL: ab_sge_i16(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<4>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b16 %rs1, [ab_sge_i16_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs2, [ab_sge_i16_param_1];
+; CHECK-NEXT:    max.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT:    cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
   %cmp = icmp sge i16 %a, %b
   %sel = select i1 %cmp, i16 %a, i16 %b
   ret i16 %sel
 }
 
 define i16 @ab_slt_i16(i16 %a, i16 %b) {
-; CHECK-LABEL: @ab_slt_i16
-; CHECK: min.s16
+; CHECK-LABEL: ab_slt_i16(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<4>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b16 %rs1, [ab_slt_i16_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs2, [ab_slt_i16_param_1];
+; CHECK-NEXT:    min.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT:    cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
   %cmp = icmp slt i16 %a, %b
   %sel = select i1 %cmp, i16 %a, i16 %b
   ret i16 %sel
 }
 
 define i16 @ab_sle_i16(i16 %a, i16 %b) {
-; CHECK-LABEL: @ab_sle_i16
-; CHECK: min.s16
+; CHECK-LABEL: ab_sle_i16(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<4>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b16 %rs1, [ab_sle_i16_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs2, [ab_sle_i16_param_1];
+; CHECK-NEXT:    min.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT:    cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
   %cmp = icmp sle i16 %a, %b
   %sel = select i1 %cmp, i16 %a, i16 %b
   ret i16 %sel
@@ -93,32 +186,72 @@ define i16 @ab_sle_i16(i16 %a, i16 %b) {
 
 ; *** ba, unsigned, i16
 define i16 @ba_ugt_i16(i16 %a, i16 %b) {
-; CHECK-LABEL: @ba_ugt_i16
-; CHECK: min.u16
+; CHECK-LABEL: ba_ugt_i16(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<4>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b16 %rs1, [ba_ugt_i16_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs2, [ba_ugt_i16_param_1];
+; CHECK-NEXT:    min.u16 %rs3, %rs1, %rs2;
+; CHECK-NEXT:    cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
   %cmp = icmp ugt i16 %a, %b
   %sel = select i1 %cmp, i16 %b, i16 %a
   ret i16 %sel
 }
 
 define i16 @ba_uge_i16(i16 %a, i16 %b) {
-; CHECK-LABEL: @ba_uge_i16
-; CHECK: min.u16
+; CHECK-LABEL: ba_uge_i16(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<4>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b16 %rs1, [ba_uge_i16_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs2, [ba_uge_i16_param_1];
+; CHECK-NEXT:    min.u16 %rs3, %rs1, %rs2;
+; CHECK-NEXT:    cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
   %cmp = icmp uge i16 %a, %b
   %sel = select i1 %cmp, i16 %b, i16 %a
   ret i16 %sel
 }
 
 define i16 @ba_ult_i16(i16 %a, i16 %b) {
-; CHECK-LABEL: @ba_ult_i16
-; CHECK: max.u16
+; CHECK-LABEL: ba_ult_i16(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<4>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b16 %rs1, [ba_ult_i16_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs2, [ba_ult_i16_param_1];
+; CHECK-NEXT:    max.u16 %rs3, %rs1, %rs2;
+; CHECK-NEXT:    cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
   %cmp = icmp ult i16 %a, %b
   %sel = select i1 %cmp, i16 %b, i16 %a
   ret i16 %sel
 }
 
 define i16 @ba_ule_i16(i16 %a, i16 %b) {
-; CHECK-LABEL: @ba_ule_i16
-; CHECK: max.u16
+; CHECK-LABEL: ba_ule_i16(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<4>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b16 %rs1, [ba_ule_i16_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs2, [ba_ule_i16_param_1];
+; CHECK-NEXT:    max.u16 %rs3, %rs1, %rs2;
+; CHECK-NEXT:    cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
   %cmp = icmp ule i16 %a, %b
   %sel = select i1 %cmp, i16 %b, i16 %a
   ret i16 %sel
@@ -126,32 +259,72 @@ define i16 @ba_ule_i16(i16 %a, i16 %b) {
 
 ; *** ba, signed, i16
 define i16 @ba_sgt_i16(i16 %a, i16 %b) {
-; CHECK-LABEL: @ba_sgt_i16
-; CHECK: min.s16
+; CHECK-LABEL: ba_sgt_i16(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<4>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b16 %rs1, [ba_sgt_i16_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs2, [ba_sgt_i16_param_1];
+; CHECK-NEXT:    min.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT:    cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
   %cmp = icmp sgt i16 %a, %b
   %sel = select i1 %cmp, i16 %b, i16 %a
   ret i16 %sel
 }
 
 define i16 @ba_sge_i16(i16 %a, i16 %b) {
-; CHECK-LABEL: @ba_sge_i16
-; CHECK: min.s16
+; CHECK-LABEL: ba_sge_i16(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<4>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b16 %rs1, [ba_sge_i16_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs2, [ba_sge_i16_param_1];
+; CHECK-NEXT:    min.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT:    cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
   %cmp = icmp sge i16 %a, %b
   %sel = select i1 %cmp, i16 %b, i16 %a
   ret i16 %sel
 }
 
 define i16 @ba_slt_i16(i16 %a, i16 %b) {
-; CHECK-LABEL: @ba_slt_i16
-; CHECK: max.s16
+; CHECK-LABEL: ba_slt_i16(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<4>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b16 %rs1, [ba_slt_i16_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs2, [ba_slt_i16_param_1];
+; CHECK-NEXT:    max.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT:    cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
   %cmp = icmp slt i16 %a, %b
   %sel = select i1 %cmp, i16 %b, i16 %a
   ret i16 %sel
 }
 
 define i16 @ba_sle_i16(i16 %a, i16 %b) {
-; CHECK-LABEL: @ba_sle_i16
-; CHECK: max.s16
+; CHECK-LABEL: ba_sle_i16(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<4>;
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b16 %rs1, [ba_sle_i16_param_0];
+; CHECK-NEXT:    ld.param.b16 %rs2, [ba_sle_i16_param_1];
+; CHECK-NEXT:    max.s16 %rs3, %rs1, %rs2;
+; CHECK-NEXT:    cvt.u32.u16 %r1, %rs3;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
+; CHECK-NEXT:    ret;
   %cmp = icmp sle i16 %a, %b
   %sel = select i1 %cmp, i16 %b, i16 %a
   ret i16 %sel
@@ -162,32 +335,64 @@ define i16 @ba_sle_i16(i16 %a, i16 %b) {
 
 ; *** ab, unsigned, i32
 define i32 @ab_ugt_i32(i32 %a, i32 %b) {
-; CHECK-LABEL: @ab_ugt_i32
-; CHECK: max.u32
+; CHECK-LABEL: ab_ugt_i32(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b32 %r1, [ab_ugt_i32_param_0];
+; CHECK-NEXT:    ld.param.b32 %r2, [ab_ugt_i32_param_1];
+; CHECK-NEXT:    max.u32 %r3, %r1, %r2;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT:    ret;
   %cmp = icmp ugt i32 %a, %b
   %sel = select i1 %cmp, i32 %a, i32 %b
   ret i32 %sel
 }
 
 define i32 @ab_uge_i32(i32 %a, i32 %b) {
-; CHECK-LABEL: @ab_uge_i32
-; CHECK: max.u32
+; CHECK-LABEL: ab_uge_i32(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b32 %r1, [ab_uge_i32_param_0];
+; CHECK-NEXT:    ld.param.b32 %r2, [ab_uge_i32_param_1];
+; CHECK-NEXT:    max.u32 %r3, %r1, %r2;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT:    ret;
   %cmp = icmp uge i32 %a, %b
   %sel = select i1 %cmp, i32 %a, i32 %b
   ret i32 %sel
 }
 
 define i32 @ab_ult_i32(i32 %a, i32 %b) {
-; CHECK-LABEL: @ab_ult_i32
-; CHECK: min.u32
+; CHECK-LABEL: ab_ult_i32(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b32 %r1, [ab_ult_i32_param_0];
+; CHECK-NEXT:    ld.param.b32 %r2, [ab_ult_i32_param_1];
+; CHECK-NEXT:    min.u32 %r3, %r1, %r2;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT:    ret;
   %cmp = icmp ult i32 %a, %b
   %sel = select i1 %cmp, i32 %a, i32 %b
   ret i32 %sel
 }
 
 define i32 @ab_ule_i32(i32 %a, i32 %b) {
-; CHECK-LABEL: @ab_ule_i32
-; CHECK: min.u32
+; CHECK-LABEL: ab_ule_i32(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b32 %r1, [ab_ule_i32_param_0];
+; CHECK-NEXT:    ld.param.b32 %r2, [ab_ule_i32_param_1];
+; CHECK-NEXT:    min.u32 %r3, %r1, %r2;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT:    ret;
   %cmp = icmp ule i32 %a, %b
   %sel = select i1 %cmp, i32 %a, i32 %b
   ret i32 %sel
@@ -195,32 +400,64 @@ define i32 @ab_ule_i32(i32 %a, i32 %b) {
 
 ; *** ab, signed, i32
 define i32 @ab_sgt_i32(i32 %a, i32 %b) {
-; CHECK-LABEL: @ab_sgt_i32
-; CHECK: max.s32
+; CHECK-LABEL: ab_sgt_i32(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b32 %r1, [ab_sgt_i32_param_0];
+; CHECK-NEXT:    ld.param.b32 %r2, [ab_sgt_i32_param_1];
+; CHECK-NEXT:    max.s32 %r3, %r1, %r2;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT:    ret;
   %cmp = icmp sgt i32 %a, %b
   %sel = select i1 %cmp, i32 %a, i32 %b
   ret i32 %sel
 }
 
 define i32 @ab_sge_i32(i32 %a, i32 %b) {
-; CHECK-LABEL: @ab_sge_i32
-; CHECK: max.s32
+; CHECK-LABEL: ab_sge_i32(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b32 %r1, [ab_sge_i32_param_0];
+; CHECK-NEXT:    ld.param.b32 %r2, [ab_sge_i32_param_1];
+; CHECK-NEXT:    max.s32 %r3, %r1, %r2;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT:    ret;
   %cmp = icmp sge i32 %a, %b
   %sel = select i1 %cmp, i32 %a, i32 %b
   ret i32 %sel
 }
 
 define i32 @ab_slt_i32(i32 %a, i32 %b) {
-; CHECK-LABEL: @ab_slt_i32
-; CHECK: min.s32
+; CHECK-LABEL: ab_slt_i32(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b32 %r1, [ab_slt_i32_param_0];
+; CHECK-NEXT:    ld.param.b32 %r2, [ab_slt_i32_param_1];
+; CHECK-NEXT:    min.s32 %r3, %r1, %r2;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT:    ret;
   %cmp = icmp slt i32 %a, %b
   %sel = select i1 %cmp, i32 %a, i32 %b
   ret i32 %sel
 }
 
 define i32 @ab_sle_i32(i32 %a, i32 %b) {
-; CHECK-LABEL: @ab_sle_i32
-; CHECK: min.s32
+; CHECK-LABEL: ab_sle_i32(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b32 %r1, [ab_sle_i32_param_0];
+; CHECK-NEXT:    ld.param.b32 %r2, [ab_sle_i32_param_1];
+; CHECK-NEXT:    min.s32 %r3, %r1, %r2;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT:    ret;
   %cmp = icmp sle i32 %a, %b
   %sel = select i1 %cmp, i32 %a, i32 %b
   ret i32 %sel
@@ -228,32 +465,64 @@ define i32 @ab_sle_i32(i32 %a, i32 %b) {
 
 ; *** ba, unsigned, i32
 define i32 @ba_ugt_i32(i32 %a, i32 %b) {
-; CHECK-LABEL: @ba_ugt_i32
-; CHECK: min.u32
+; CHECK-LABEL: ba_ugt_i32(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b32 %r1, [ba_ugt_i32_param_0];
+; CHECK-NEXT:    ld.param.b32 %r2, [ba_ugt_i32_param_1];
+; CHECK-NEXT:    min.u32 %r3, %r1, %r2;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT:    ret;
   %cmp = icmp ugt i32 %a, %b
   %sel = select i1 %cmp, i32 %b, i32 %a
   ret i32 %sel
 }
 
 define i32 @ba_uge_i32(i32 %a, i32 %b) {
-; CHECK-LABEL: @ba_uge_i32
-; CHECK: min.u32
+; CHECK-LABEL: ba_uge_i32(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.b32 %r1, [ba_uge_i32_param_0];
+; CHECK-NEXT:    ld.param.b32 %r2, [ba_uge_i32_param_1];
+; CHECK-NEXT:    min.u32 %r3, %r1, %r2;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT:    ret;
   %cmp = icmp uge i32 %a, %b
   %sel = select i1 %cmp, i32 %b, i32 %a
   ret i32 %sel
 }
 
 define i32 @ba_ult_i32(i32 %a, i32 %b) {
-; CHECK-LABEL: @ba_ult_i32
-; CHECK: max.u32
+; CHECK-LABEL: ba_ult_i32(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<4>;
+; CHECK-EMPTY:
+; CHE...
[truncated]

Artem-B

LGTM in principle, modulo the question of whether we want to make lowering depend on optimization level.

Artem-B · 2025-08-04T22:11:14Z

llvm/lib/Target/NVPTX/NVPTXInstrInfo.td

 def SMIN16x2 : I16x2<"min.s", smin>;
 def UMIN16x2 : I16x2<"min.u", umin>;

+let Predicates = [hasPTX<80>, hasSM<90>, hasOptEnabled] in {


Is there a particular reason to keep lowering to max.relu behind optimization level?
We virtually never do that for other patterns we match to an instruction.

The way I see it, the job of the lowering is to always convert the DAG to the best instructions available to represent it. User-controlled optimizations generally belong to the higher levels. I would prefer not to add another moving piece into the lowering process.

AFAICT TM.getOptLevel() is mostly used to configure which passes get enabled/disabled, and NVPTX is the only back-end which uses it in the back-end tablegen file.

There is not a particular reason. I suppose that another perspective would be that at -O0 we want to be able to turn off as much as possible for debugging purposes. I of course don't have any reason to think there is anything wrong with these folds but it might help debugging any issues if it were possible to turn optional folds off in general. In any case, these seem pretty innocuous so I'll go ahead and remove this check.

Artem-B · 2025-08-04T22:15:05Z

llvm/test/CodeGen/NVPTX/combine-min-max.ll

+; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -O3 | FileCheck %s
+; RUN: %if ptxas-12.0 %{ llc < %s -mcpu=sm_90 -mattr=+ptx80 -O3 | %ptxas-verify %}


We may want to keep pre-sm_90 runs, too.

llvm-ci · 2025-08-05T04:12:01Z

LLVM Buildbot has detected a new failure on builder llvm-nvptx-nvidia-win running on as-builder-8 while building llvm at step 7 "test-build-unified-tree-check-llvm".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/54/builds/11525

Here is the relevant piece of the build log for the reference

Step 7 (test-build-unified-tree-check-llvm) failure: test (failure)
******************** TEST 'LLVM :: CodeGen/NVPTX/combine-min-max.ll' FAILED ********************
Exit Code: 4294967295

Command Output (stdout):
--
# RUN: at line 2
c:\buildbot\as-builder-8\llvm-nvptx-nvidia-win\build\bin\llc.exe < C:\buildbot\as-builder-8\llvm-nvptx-nvidia-win\llvm-project\llvm\test\CodeGen\NVPTX\combine-min-max.ll -mcpu=sm_90 -mattr=+ptx80 -O3 | c:\buildbot\as-builder-8\llvm-nvptx-nvidia-win\build\bin\filecheck.exe C:\buildbot\as-builder-8\llvm-nvptx-nvidia-win\llvm-project\llvm\test\CodeGen\NVPTX\combine-min-max.ll --check-prefixes=CHECK,SM90
# executed command: 'c:\buildbot\as-builder-8\llvm-nvptx-nvidia-win\build\bin\llc.exe' -mcpu=sm_90 -mattr=+ptx80 -O3
# executed command: 'c:\buildbot\as-builder-8\llvm-nvptx-nvidia-win\build\bin\filecheck.exe' 'C:\buildbot\as-builder-8\llvm-nvptx-nvidia-win\llvm-project\llvm\test\CodeGen\NVPTX\combine-min-max.ll' --check-prefixes=CHECK,SM90
# RUN: at line 3
c:\buildbot\as-builder-8\llvm-nvptx-nvidia-win\build\bin\llc.exe < C:\buildbot\as-builder-8\llvm-nvptx-nvidia-win\llvm-project\llvm\test\CodeGen\NVPTX\combine-min-max.ll -mcpu=sm_20 -O3 | c:\buildbot\as-builder-8\llvm-nvptx-nvidia-win\build\bin\filecheck.exe C:\buildbot\as-builder-8\llvm-nvptx-nvidia-win\llvm-project\llvm\test\CodeGen\NVPTX\combine-min-max.ll --check-prefixes=CHECK,SM20
# executed command: 'c:\buildbot\as-builder-8\llvm-nvptx-nvidia-win\build\bin\llc.exe' -mcpu=sm_20 -O3
# executed command: 'c:\buildbot\as-builder-8\llvm-nvptx-nvidia-win\build\bin\filecheck.exe' 'C:\buildbot\as-builder-8\llvm-nvptx-nvidia-win\llvm-project\llvm\test\CodeGen\NVPTX\combine-min-max.ll' --check-prefixes=CHECK,SM20
# RUN: at line 4
c:\buildbot\as-builder-8\llvm-nvptx-nvidia-win\build\bin\llc.exe < C:\buildbot\as-builder-8\llvm-nvptx-nvidia-win\llvm-project\llvm\test\CodeGen\NVPTX\combine-min-max.ll -mcpu=sm_90 -mattr=+ptx80 -O3 | c:/buildbot/latest-cuda/bin/ptxas.exe -arch=sm_60 -c -
# executed command: 'c:\buildbot\as-builder-8\llvm-nvptx-nvidia-win\build\bin\llc.exe' -mcpu=sm_90 -mattr=+ptx80 -O3
# executed command: c:/buildbot/latest-cuda/bin/ptxas.exe -arch=sm_60 -c -
# .---command stdout------------
# | ptxas fatal   : SM version specified by .target is higher than default SM version assumed
# `-----------------------------
# error: command failed with exit status: 0xffffffff

--

********************

llvm-ci · 2025-08-05T04:25:11Z

LLVM Buildbot has detected a new failure on builder llvm-nvptx64-nvidia-ubuntu running on as-builder-7 while building llvm at step 6 "test-build-unified-tree-check-llvm".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/160/builds/22440

Here is the relevant piece of the build log for the reference

Step 6 (test-build-unified-tree-check-llvm) failure: test (failure)
******************** TEST 'LLVM :: CodeGen/NVPTX/combine-min-max.ll' FAILED ********************
Exit Code: 255

Command Output (stderr):
--
/home/buildbot/worker/as-builder-7/llvm-nvptx64-nvidia-ubuntu/build/bin/llc < /home/buildbot/worker/as-builder-7/llvm-nvptx64-nvidia-ubuntu/llvm-project/llvm/test/CodeGen/NVPTX/combine-min-max.ll -mcpu=sm_90 -mattr=+ptx80 -O3 | /home/buildbot/worker/as-builder-7/llvm-nvptx64-nvidia-ubuntu/build/bin/FileCheck /home/buildbot/worker/as-builder-7/llvm-nvptx64-nvidia-ubuntu/llvm-project/llvm/test/CodeGen/NVPTX/combine-min-max.ll --check-prefixes=CHECK,SM90 # RUN: at line 2
+ /home/buildbot/worker/as-builder-7/llvm-nvptx64-nvidia-ubuntu/build/bin/llc -mcpu=sm_90 -mattr=+ptx80 -O3
+ /home/buildbot/worker/as-builder-7/llvm-nvptx64-nvidia-ubuntu/build/bin/FileCheck /home/buildbot/worker/as-builder-7/llvm-nvptx64-nvidia-ubuntu/llvm-project/llvm/test/CodeGen/NVPTX/combine-min-max.ll --check-prefixes=CHECK,SM90
/home/buildbot/worker/as-builder-7/llvm-nvptx64-nvidia-ubuntu/build/bin/llc < /home/buildbot/worker/as-builder-7/llvm-nvptx64-nvidia-ubuntu/llvm-project/llvm/test/CodeGen/NVPTX/combine-min-max.ll -mcpu=sm_20 -O3 | /home/buildbot/worker/as-builder-7/llvm-nvptx64-nvidia-ubuntu/build/bin/FileCheck /home/buildbot/worker/as-builder-7/llvm-nvptx64-nvidia-ubuntu/llvm-project/llvm/test/CodeGen/NVPTX/combine-min-max.ll --check-prefixes=CHECK,SM20 # RUN: at line 3
+ /home/buildbot/worker/as-builder-7/llvm-nvptx64-nvidia-ubuntu/build/bin/FileCheck /home/buildbot/worker/as-builder-7/llvm-nvptx64-nvidia-ubuntu/llvm-project/llvm/test/CodeGen/NVPTX/combine-min-max.ll --check-prefixes=CHECK,SM20
+ /home/buildbot/worker/as-builder-7/llvm-nvptx64-nvidia-ubuntu/build/bin/llc -mcpu=sm_20 -O3
/home/buildbot/worker/as-builder-7/llvm-nvptx64-nvidia-ubuntu/build/bin/llc < /home/buildbot/worker/as-builder-7/llvm-nvptx64-nvidia-ubuntu/llvm-project/llvm/test/CodeGen/NVPTX/combine-min-max.ll -mcpu=sm_90 -mattr=+ptx80 -O3 | /usr/local/cuda/bin/ptxas -arch=sm_60 -c - # RUN: at line 4
+ /home/buildbot/worker/as-builder-7/llvm-nvptx64-nvidia-ubuntu/build/bin/llc -mcpu=sm_90 -mattr=+ptx80 -O3
+ /usr/local/cuda/bin/ptxas -arch=sm_60 -c -
ptxas fatal   : SM version specified by .target is higher than default SM version assumed

--

********************

llvm-ci · 2025-08-05T04:29:28Z

LLVM Buildbot has detected a new failure on builder llvm-nvptx-nvidia-ubuntu running on as-builder-7 while building llvm at step 6 "test-build-unified-tree-check-llvm".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/180/builds/22583

Here is the relevant piece of the build log for the reference

Step 6 (test-build-unified-tree-check-llvm) failure: test (failure)
******************** TEST 'LLVM :: CodeGen/NVPTX/combine-min-max.ll' FAILED ********************
Exit Code: 255

Command Output (stderr):
--
/home/buildbot/worker/as-builder-7/llvm-nvptx-nvidia-ubuntu/build/bin/llc < /home/buildbot/worker/as-builder-7/llvm-nvptx-nvidia-ubuntu/llvm-project/llvm/test/CodeGen/NVPTX/combine-min-max.ll -mcpu=sm_90 -mattr=+ptx80 -O3 | /home/buildbot/worker/as-builder-7/llvm-nvptx-nvidia-ubuntu/build/bin/FileCheck /home/buildbot/worker/as-builder-7/llvm-nvptx-nvidia-ubuntu/llvm-project/llvm/test/CodeGen/NVPTX/combine-min-max.ll --check-prefixes=CHECK,SM90 # RUN: at line 2
+ /home/buildbot/worker/as-builder-7/llvm-nvptx-nvidia-ubuntu/build/bin/llc -mcpu=sm_90 -mattr=+ptx80 -O3
+ /home/buildbot/worker/as-builder-7/llvm-nvptx-nvidia-ubuntu/build/bin/FileCheck /home/buildbot/worker/as-builder-7/llvm-nvptx-nvidia-ubuntu/llvm-project/llvm/test/CodeGen/NVPTX/combine-min-max.ll --check-prefixes=CHECK,SM90
/home/buildbot/worker/as-builder-7/llvm-nvptx-nvidia-ubuntu/build/bin/llc < /home/buildbot/worker/as-builder-7/llvm-nvptx-nvidia-ubuntu/llvm-project/llvm/test/CodeGen/NVPTX/combine-min-max.ll -mcpu=sm_20 -O3 | /home/buildbot/worker/as-builder-7/llvm-nvptx-nvidia-ubuntu/build/bin/FileCheck /home/buildbot/worker/as-builder-7/llvm-nvptx-nvidia-ubuntu/llvm-project/llvm/test/CodeGen/NVPTX/combine-min-max.ll --check-prefixes=CHECK,SM20 # RUN: at line 3
+ /home/buildbot/worker/as-builder-7/llvm-nvptx-nvidia-ubuntu/build/bin/llc -mcpu=sm_20 -O3
+ /home/buildbot/worker/as-builder-7/llvm-nvptx-nvidia-ubuntu/build/bin/FileCheck /home/buildbot/worker/as-builder-7/llvm-nvptx-nvidia-ubuntu/llvm-project/llvm/test/CodeGen/NVPTX/combine-min-max.ll --check-prefixes=CHECK,SM20
/home/buildbot/worker/as-builder-7/llvm-nvptx-nvidia-ubuntu/build/bin/llc < /home/buildbot/worker/as-builder-7/llvm-nvptx-nvidia-ubuntu/llvm-project/llvm/test/CodeGen/NVPTX/combine-min-max.ll -mcpu=sm_90 -mattr=+ptx80 -O3 | /usr/local/cuda/bin/ptxas -arch=sm_60 -c - # RUN: at line 4
+ /home/buildbot/worker/as-builder-7/llvm-nvptx-nvidia-ubuntu/build/bin/llc -mcpu=sm_90 -mattr=+ptx80 -O3
+ /usr/local/cuda/bin/ptxas -arch=sm_60 -c -
ptxas fatal   : SM version specified by .target is higher than default SM version assumed

--

********************

[NVPTX] Add support for integer min/max ReLU idiom

0995591

AlexMaclean requested review from Artem-B and justinfargnoli August 1, 2025 16:40

AlexMaclean self-assigned this Aug 1, 2025

llvmbot added the backend:NVPTX label Aug 1, 2025

justinfargnoli approved these changes Aug 1, 2025

View reviewed changes

Artem-B approved these changes Aug 4, 2025

View reviewed changes

AlexMaclean added 2 commits August 4, 2025 22:53

address comments

f852232

Merge branch 'main' into dev/amaclean/upstream/min-max-relu

b2aec27

AlexMaclean merged commit 5678aef into llvm:main Aug 5, 2025
9 checks passed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Uh oh!

[NVPTX] Add support for integer min/max ReLU idiom #151727

[NVPTX] Add support for integer min/max ReLU idiom #151727

Uh oh!

AlexMaclean commented Aug 1, 2025

Uh oh!

llvmbot commented Aug 1, 2025

Uh oh!

Artem-B left a comment

Uh oh!

Artem-B Aug 4, 2025

Uh oh!

AlexMaclean Aug 4, 2025

Uh oh!

Artem-B Aug 4, 2025

Uh oh!

AlexMaclean Aug 5, 2025

Uh oh!

Uh oh!

llvm-ci commented Aug 5, 2025

Uh oh!

llvm-ci commented Aug 5, 2025

Uh oh!

llvm-ci commented Aug 5, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

		; RUN: llc < %s -mcpu=sm_90 -mattr=+ptx80 -O3 \| FileCheck %s
		; RUN: %if ptxas-12.0 %{ llc < %s -mcpu=sm_90 -mattr=+ptx80 -O3 \| %ptxas-verify %}

[NVPTX] Add support for integer min/max ReLU idiom #151727

[NVPTX] Add support for integer min/max ReLU idiom #151727

Uh oh!

Conversation

AlexMaclean commented Aug 1, 2025

Uh oh!

llvmbot commented Aug 1, 2025

Uh oh!

Artem-B left a comment

Choose a reason for hiding this comment

Uh oh!

Artem-B Aug 4, 2025

Choose a reason for hiding this comment

Uh oh!

AlexMaclean Aug 4, 2025

Choose a reason for hiding this comment

Uh oh!

Artem-B Aug 4, 2025

Choose a reason for hiding this comment

Uh oh!

AlexMaclean Aug 5, 2025

Choose a reason for hiding this comment

Uh oh!

Uh oh!

llvm-ci commented Aug 5, 2025

Uh oh!

llvm-ci commented Aug 5, 2025

Uh oh!

llvm-ci commented Aug 5, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants