From 12a4dfe833723b7b6c4865b91667cd6b310417b5 Mon Sep 17 00:00:00 2001
From: Aakanksha Patil <aakanksha555@gmail.com>
Date: Wed, 12 Aug 2020 13:51:31 -0400
Subject: [PATCH 1/2] Add gfx1031

Change-Id: Ie12fd9ffe7d7b51fadfe67029c2271ec6779828c
---
 oclc/src/isa_version_1031.cl | 11 +++++++++++
 1 file changed, 11 insertions(+)
 create mode 100644 oclc/src/isa_version_1031.cl

diff --git a/oclc/src/isa_version_1031.cl b/oclc/src/isa_version_1031.cl
new file mode 100644
index 0000000..4dc1b88
--- /dev/null
+++ b/oclc/src/isa_version_1031.cl
@@ -0,0 +1,11 @@
+/*===--------------------------------------------------------------------------
+ *                   ROCm Device Libraries
+ *
+ * This file is distributed under the University of Illinois Open Source
+ * License. See LICENSE.TXT for details.
+ *===------------------------------------------------------------------------*/
+
+#include "oclc.h"
+
+const __constant int __oclc_ISA_version = 10301;
+

From 41b390cd90f7f84605997657c31eb0d8933ad64a Mon Sep 17 00:00:00 2001
From: Brian Sumner <brian.sumner@amd.com>
Date: Tue, 18 Aug 2020 10:36:14 -0700
Subject: [PATCH 2/2] Switch a few intrinsics to builtins and add required
 features

Change-Id: Ie6b716946258a066d101e64e1a6099fa4bf1bb4b
---
 ockl/src/wfredscan.cl | 317 ++++++++++++++++++++++++------------------
 1 file changed, 184 insertions(+), 133 deletions(-)

diff --git a/ockl/src/wfredscan.cl b/ockl/src/wfredscan.cl
index c11c47b..8e99037 100644
--- a/ockl/src/wfredscan.cl
+++ b/ockl/src/wfredscan.cl
@@ -60,7 +60,7 @@
 #define half_swizzle(X,Y) AS_HALF((ushort)uint_swizzle((uint)AS_USHORT(X),Y))
 
 // DPP16
-#define uint_dpp(ID,X,C,R,B,W) __llvm_amdgcn_update_dpp_i32(ID,X,C,R,B,W)
+#define uint_dpp(ID,X,C,R,B,W) __builtin_amdgcn_update_dpp(ID,X,C,R,B,W)
 #define ulong_dpp(ID,X,C,R,B,W) ({ \
     uint2 __x = AS_UINT2(X); \
     uint2 __r; \
@@ -68,14 +68,14 @@
     __r.hi = uint_dpp((uint)(ID >> 32), __x.hi, C, R, B, W); \
     AS_ULONG(__r); \
 })
-#define int_dpp(ID,X,C,R,B,W) (int)uint_dpp((uint)ID,X,C,R,B,W)
-#define long_dpp(ID,X,C,R,B,W) (long)ulong_dpp((ulong)ID,(ulong)X,C,R,B,W)
+#define int_dpp(ID,X,C,R,B,W) AS_INT(uint_dpp(AS_UINT(ID),AS_UINT(X),C,R,B,W))
+#define long_dpp(ID,X,C,R,B,W) AS_LONG(ulong_dpp(AS_ULONG(ID),AS_ULONG(X),C,R,B,W))
 #define float_dpp(ID,X,C,R,B,W) AS_FLOAT(uint_dpp(AS_UINT(ID),AS_UINT(X),C,R,B,W))
 #define double_dpp(ID,X,C,R,B,W) AS_DOUBLE(ulong_dpp(AS_ULONG(ID),AS_ULONG(X),C,R,B,W))
 #define half_dpp(ID,X,C,R,B,W) AS_HALF((ushort)uint_dpp((uint)AS_USHORT(ID),(uint)AS_USHORT(X),C,R,B,W))
 
 // DPP8
-#define uint_dpp8(X,S) __llvm_amdgcn_mov_dpp8_i32(X,S)
+#define uint_dpp8(X,S) __builtin_amdgcn_mov_dpp8(X,S)
 #define ulong_dpp8(X,S) ({ \
     uint2 __x = AS_UINT2(X); \
     uint2 __r; \
@@ -90,7 +90,7 @@
 #define half_dpp8(X,S) AS_HALF((ushort)uint_dpp8((uint)AS_USHORT(X),S))
 
 // permlane16
-#define uint_permlane16(ID,X,S0,S1,W) __llvm_amdgcn_permlane16(ID,X,S0,S1,false,W)
+#define uint_permlane16(ID,X,S0,S1,W) __builtin_amdgcn_permlane16(ID,X,S0,S1,false,W)
 #define ulong_permlane16(ID,X,S0,S1,W) ({ \
     uint2 __x = AS_UINT2(X); \
     uint2 __r; \
@@ -98,14 +98,14 @@
     __r.hi = uint_permlane16((uint)(ID>>32),__x.hi,S0,S1,W); \
     AS_ULONG(__r); \
 })
-#define int_permlane16(ID,X,S0,S1,W) (int)uint_permlane16((uint)ID,(uint)X,S0,S1,W)
-#define long_permlane16(ID,X,S0,S1,W) (long)ulong_permlane16((ulong)ID,(ulong)X,S0,S1,W)
+#define int_permlane16(ID,X,S0,S1,W) AS_INT(uint_permlane16(AS_UINT(ID),AS_UINT(X),S0,S1,W))
+#define long_permlane16(ID,X,S0,S1,W) AS_LONG(ulong_permlane16(AS_ULONG(ID),AS_ULONG(X),S0,S1,W))
 #define float_permlane16(ID, X,S0,S1,W) AS_FLOAT(uint_permlane16(AS_UINT(ID),AS_UINT(X),S0,S1,W))
 #define double_permlane16(ID, X,S0,S1,W) AS_DOUBLE(ulong_permlane16(AS_ULONG(ID),AS_ULONG(X),S0,S1,W))
 #define half_permlane16(ID,X,S0,S1,W) AS_HALF((ushort)uint_permlane16((uint)AS_USHORT(ID),(uint)AS_USHORT(X),S0,S1,W))
 
 // permlanex16
-#define uint_permlanex16(ID,X,S0,S1,W) __llvm_amdgcn_permlanex16(ID,X,S0,S1,false,W)
+#define uint_permlanex16(ID,X,S0,S1,W) __builtin_amdgcn_permlanex16(ID,X,S0,S1,false,W)
 #define ulong_permlanex16(ID,X,S0,S1,W) ({ \
     uint2 __x = AS_UINT2(X); \
     uint2 __r; \
@@ -113,8 +113,8 @@
     __r.hi = uint_permlanex16((uint)(ID>>32),__x.hi,S0,S1,W); \
     AS_ULONG(__r); \
 })
-#define int_permlanex16(ID,X,S0,S1,W) (int)uint_permlanex16((uint)ID,(uint)X,S0,S1,W)
-#define long_permlanex16(ID,X,S0,S1,W) (long)ulong_permlanex16((ulong)ID,(ulong)X,S0,S1,W)
+#define int_permlanex16(ID,X,S0,S1,W) AS_INT(uint_permlanex16(AS_UINT(ID),AS_UINT(X),S0,S1,W))
+#define long_permlanex16(ID,X,S0,S1,W) AS_LONG(ulong_permlanex16(AS_ULONG(ID),AS_ULONG(X),S0,S1,W))
 #define float_permlanex16(ID, X,S0,S1,W) AS_FLOAT(uint_permlanex16(AS_UINT(ID),AS_UINT(X),S0,S1,W))
 #define double_permlanex16(ID, X,S0,S1,W) AS_DOUBLE(ulong_permlanex16(AS_ULONG(ID),AS_ULONG(X),S0,S1,W))
 #define half_permlanex16(ID,X,S0,S1,W) AS_HALF((ushort)uint_permlanex16((uint)AS_USHORT(ID),(uint)AS_USHORT(X),S0,S1,W))
@@ -211,10 +211,11 @@ GENMAX(ulong)
 #define long_xor(X,Y) XOR(X,Y)
 
 
-// Reduce with operation OP over full wave using swizzle
-// Input in x, r is result
-#define RED_GFX7_FULL(T,OP) \
-    T v; \
+#define GENRED7_FULL(T,OP,ID,IDZ) \
+static T \
+red7_full_##T##_##OP(T x) \
+{ \
+    T v, r; \
  \
     v = T##_swizzle(x, SWIZZLE_QUAD_PERM(0x1,0x0,0x3,0x2)); \
     r = T##_##OP(x, v); \
@@ -231,12 +232,17 @@ GENMAX(ulong)
     v = T##_swizzle(r, SWIZZLE_32_LIMITED(0x1f,0x00,0x10)); \
     r = T##_##OP(r, v); \
  \
-    r = T##_##OP(T##_readlane(r,0), T##_readlane(r,32))
+    r = T##_##OP(T##_readlane(r,0), T##_readlane(r,32)); \
+ \
+    return r; \
+}
 
-// Reduce with operation OP over partial wave using swizzle
-// Input in x, r is result
-#define RED_GFX7_PART(T,OP,ID) \
-    if (ID == (T)0) { \
+#define GENRED7_PART(T,OP,ID,IDZ) \
+static T \
+red7_part_##T##_##OP(T x) \
+{ \
+    T r; \
+    if (IDZ) { \
         T v; \
  \
         v = T##_swizzle(x,    SWIZZLE_QUAD_PERM(0x1,0x0,0x3,0x2)); \
@@ -289,64 +295,79 @@ GENMAX(ulong)
         t = T##_readlane(r, 32); \
         v = (__builtin_amdgcn_read_exec_hi() & 1) ? t : ID; \
         r = T##_##OP(T##_readlane(r, 0), v); \
-    }
+    } \
+ \
+    return r; \
+}
 
+#define GENRED7(T,OP,ID,IDZ) \
+    GENRED7_FULL(T,OP,ID,IDZ) \
+    GENRED7_PART(T,OP,ID,IDZ)
 
-// Reduce with operation OP using DPP
-// Input in x, r is result
-#define RED_GFX89(T,OP,ID) \
-    T v; \
+#define GENRED89(T,OP,ID,IDZ) \
+__attribute__((target("dpp"))) static T \
+red89_##T##_##OP(T x) \
+{ \
+    T r, v; \
  \
-    v = T##_dpp(ID, x, DPP_ROW_SL(1), 0xf, 0xf, ID == (T)0); \
+    v = T##_dpp(ID, x, DPP_ROW_SL(1), 0xf, 0xf, IDZ); \
     r = T##_##OP(x, v); \
  \
-    v = T##_dpp(ID, r, DPP_ROW_SL(2), 0xf, 0xf, ID == (T)0); \
+    v = T##_dpp(ID, r, DPP_ROW_SL(2), 0xf, 0xf, IDZ); \
     r = T##_##OP(r, v); \
  \
-    v = T##_dpp(ID, r, DPP_ROW_SL(4), 0xf, 0xf, ID == (T)0); \
+    v = T##_dpp(ID, r, DPP_ROW_SL(4), 0xf, 0xf, IDZ); \
     r = T##_##OP(r, v); \
  \
-    v = T##_dpp(ID, r, DPP_ROW_SL(8), 0xf, 0xf, ID == (T)0); \
+    v = T##_dpp(ID, r, DPP_ROW_SL(8), 0xf, 0xf, IDZ); \
     r = T##_##OP(r, v); \
  \
-    v = T##_dpp(ID, r, DPP_WF_SL1, 0xf, 0xf, ID == (T)0); \
-    v = T##_dpp(ID, v, DPP_ROW_MIRROR, 0xf, 0xf, ID == (T)0); \
+    v = T##_dpp(ID, r, DPP_WF_SL1, 0xf, 0xf, IDZ); \
+    v = T##_dpp(ID, v, DPP_ROW_MIRROR, 0xf, 0xf, IDZ); \
     r = T##_##OP(r, v); \
  \
     v = T##_readlane(r, 32); \
     v = (__builtin_amdgcn_read_exec_hi() & 1) ? v : ID; \
-    r = T##_##OP(T##_readlane(r, 0), v);
+    r = T##_##OP(T##_readlane(r, 0), v); \
+ \
+    return r; \
+}
 
-// Reduce with operation OP using DPP
-// Input in x, r is result
-#define RED_GFX10(T,OP,ID) \
-    T v; \
+#define GENRED10(T,OP,ID,IDZ) \
+__attribute__((target("dpp,gfx10-insts"))) static T \
+red10_##T##_##OP(T x) \
+{ \
+    T r, v; \
  \
-    v = T##_dpp(ID, x, DPP_ROW_XMASK(0x1), 0xf, 0xf, ID == (T)0); \
+    v = T##_dpp(ID, x, DPP_ROW_XMASK(0x1), 0xf, 0xf, IDZ); \
     r = T##_##OP(x, v); \
  \
-    v = T##_dpp(ID, r, DPP_ROW_XMASK(0x2), 0xf, 0xf, ID == (T)0); \
+    v = T##_dpp(ID, r, DPP_ROW_XMASK(0x2), 0xf, 0xf, IDZ); \
     r = T##_##OP(r, v); \
  \
-    v = T##_dpp(ID, r, DPP_ROW_XMASK(0x4), 0xf, 0xf, ID == (T)0); \
+    v = T##_dpp(ID, r, DPP_ROW_XMASK(0x4), 0xf, 0xf, IDZ); \
     r = T##_##OP(r, v); \
  \
-    v = T##_dpp(ID, r, DPP_ROW_XMASK(0x8), 0xf, 0xf, ID == (T)0); \
+    v = T##_dpp(ID, r, DPP_ROW_XMASK(0x8), 0xf, 0xf, IDZ); \
     r = T##_##OP(r, v); \
  \
-    v = T##_permlanex16(ID, r, 0, 0, ID == (T)0); \
+    v = T##_permlanex16(ID, r, 0, 0, IDZ); \
     r = T##_##OP(r, v); \
  \
     if (__oclc_wavefrontsize64) { \
         T v = T##_readlane(r, 32); \
         v = (__builtin_amdgcn_read_exec_hi() & 1) ? v : ID; \
         r =  T##_##OP(T##_readlane(r, 0), v); \
-    }
+    } \
+ \
+    return r; \
+}
 
-// Inclusive scan with operation OP using swizzle
-// Input is x, l is lane, output is s
-#define ISCAN_GFX7(T,OP,ID) \
-    T v; \
+#define GENISCAN7(T,OP,ID,IDZ) \
+static T \
+iscan7_##T##_##OP(T x, uint l) \
+{ \
+    T s, v; \
  \
     v = T##_swizzle(x, SWIZZLE_32_LIMITED(0x1e,0x00,0x00)); \
     v = (l & 1) ? v : ID; \
@@ -370,24 +391,27 @@ GENMAX(ulong)
  \
     v = T##_readlane(s, 31); \
     v = l > 31 ? v : ID; \
-    s = T##_##OP(s, v)
-
+    s = T##_##OP(s, v); \
+ \
+    return s; \
+}
 
-// Inclusive scan with operation OP using DPP
-// Input is x, l is lane, output is s
-#define ISCAN_GFX89(T,OP,ID) \
-    T v; \
+#define GENISCAN89(T,OP,ID,IDZ) \
+__attribute__((target("dpp"))) static T \
+iscan89_##T##_##OP(T x, uint l) \
+{ \
+    T s, v; \
  \
-    v = T##_dpp(ID, x, DPP_ROW_SR(1), 0xf, 0xf, ID == (T)0); \
+    v = T##_dpp(ID, x, DPP_ROW_SR(1), 0xf, 0xf, IDZ); \
     s = T##_##OP(x, v); \
  \
-    v = T##_dpp(ID, s, DPP_ROW_SR(2), 0xf, 0xf, ID == (T)0); \
+    v = T##_dpp(ID, s, DPP_ROW_SR(2), 0xf, 0xf, IDZ); \
     s = T##_##OP(s, v); \
  \
-    v = T##_dpp(ID, s, DPP_ROW_SR(4), 0xf, 0xf, ID == (T)0); \
+    v = T##_dpp(ID, s, DPP_ROW_SR(4), 0xf, 0xf, IDZ); \
     s = T##_##OP(s, v); \
  \
-    v = T##_dpp(ID, s, DPP_ROW_SR(8), 0xf, 0xf, ID == (T)0); \
+    v = T##_dpp(ID, s, DPP_ROW_SR(8), 0xf, 0xf, IDZ); \
     s = T##_##OP(s, v); \
  \
     v = T##_dpp(ID, s, DPP_ROW_BCAST15, 0xa, 0xf, false); \
@@ -395,25 +419,29 @@ GENMAX(ulong)
  \
     v = T##_dpp(ID, s, DPP_ROW_BCAST31, 0xc, 0xf, false); \
     s = T##_##OP(s, v); \
+ \
+    return s; \
+}
 
-// Inclusive scan with operation OP using DPP
-// Input is x, l is lane, output is s
-#define ISCAN_GFX10(T,OP,ID) \
-    T v; \
+#define GENISCAN10(T,OP,ID,IDZ) \
+__attribute__((target("dpp,gfx10-insts"))) static T \
+iscan10_##T##_##OP(T x, uint l) \
+{ \
+    T s, v; \
  \
-    v = T##_dpp(ID, x, DPP_ROW_SR(1), 0xf, 0xf, ID == (T)0); \
+    v = T##_dpp(ID, x, DPP_ROW_SR(1), 0xf, 0xf, IDZ); \
     s = T##_##OP(x, v); \
  \
-    v = T##_dpp(ID, s, DPP_ROW_SR(2), 0xf, 0xf, ID == (T)0); \
+    v = T##_dpp(ID, s, DPP_ROW_SR(2), 0xf, 0xf, IDZ); \
     s = T##_##OP(s, v); \
  \
-    v = T##_dpp(ID, s, DPP_ROW_SR(4), 0xf, 0xf, ID == (T)0); \
+    v = T##_dpp(ID, s, DPP_ROW_SR(4), 0xf, 0xf, IDZ); \
     s = T##_##OP(s, v); \
  \
-    v = T##_dpp(ID, s, DPP_ROW_SR(8), 0xf, 0xf, ID == (T)0); \
+    v = T##_dpp(ID, s, DPP_ROW_SR(8), 0xf, 0xf, IDZ); \
     s = T##_##OP(s, v); \
  \
-    v = T##_permlanex16(ID, s, 0xffffffff, 0xffffffff, ID == (T)0); \
+    v = T##_permlanex16(ID, s, 0xffffffff, 0xffffffff, IDZ); \
     v = (l & 0x10) ? v : ID; \
     s = T##_##OP(s, v); \
  \
@@ -421,11 +449,15 @@ GENMAX(ulong)
         v = T##_readlane(s, 31); \
         v = l > 31 ? v : ID; \
         s = T##_##OP(s, v); \
-    }
+    } \
+ \
+     return s; \
+}
 
-// Shift right 1 on entire wavefront using swizzle
-// input is s, l is lane, output is s
-#define SR1_SWIZZLE(T,ID) \
+#define GENSR1_7(T,OP,ID,IDZ) \
+static T \
+sr1_7_##T##_##OP(T s, uint l) \
+{ \
     T v; \
     T t = s; \
  \
@@ -443,25 +475,35 @@ GENMAX(ulong)
     v = T##_readlane(t, 31); \
     s = l == 32 ? v : s; \
  \
-    s = l == 0 ? ID : s
+    s = l == 0 ? ID : s; \
+ \
+    return s; \
+}
 
-// Shift right 1 on entire wavefront using DPP
-// input is s, l is lane, output is s
-#define SR1_GFX89(T,ID) \
-    s = T##_dpp(ID, s, DPP_WF_SR1, 0xf, 0xf, ID == (T)0); \
 
-// Shift right 1 on entire wavefront using DPP
-// input is s, l is lane, output is s
-#define SR1_GFX10(T,ID) \
-    T t = T##_dpp(ID, s, DPP_ROW_SR(1), 0xf, 0xf, ID == (T)0); \
-    T v = T##_permlanex16(ID, s, 0xffffffff, 0xffffffff, ID == (T)0); \
+#define GENSR1_89(T,OP,ID,IDZ) \
+__attribute__((target("dpp"))) static T \
+sr1_89_##T##_##OP(T s, uint l) \
+{ \
+    return T##_dpp(ID, s, DPP_WF_SR1, 0xf, 0xf, IDZ); \
+}
+
+#define GENSR1_10(T,OP,ID,IDZ) \
+__attribute((target("dpp,gfx10-insts"))) static T \
+sr1_10_##T##_##OP(T s, uint l) \
+{ \
+    T t = T##_dpp(ID, s, DPP_ROW_SR(1), 0xf, 0xf, IDZ); \
+    T v = T##_permlanex16(ID, s, 0xffffffff, 0xffffffff, IDZ); \
     if (__oclc_wavefrontsize64) { \
         T w = T##_readlane(s, 31); \
         v = l == 32 ? w : v; \
         s = ((l == 32) | ((l & 0x1f) == 0x10)) ? v : t; \
     } else {\
         s = l == 16 ? v : t; \
-    }
+    } \
+ \
+    return s; \
+}
 
 IATTR static bool
 fullwave(void)
@@ -473,26 +515,35 @@ fullwave(void)
     }
 }
 
-#define GENRED(T,OP,ID) \
+#define GENRED(T,OP,ID,IDZ) \
+GENRED7(T,OP,ID,IDZ) \
+GENRED89(T,OP,ID,IDZ) \
+GENRED10(T,OP,ID,IDZ) \
 IATTR T \
 C(__ockl_wfred_,C(OP,T##_suf))(T x) \
 { \
     T r; \
     if (__oclc_ISA_version < 8000) { \
          if (fullwave()) { \
-             RED_GFX7_FULL(T,OP); \
+             r = red7_full_##T##_##OP(x); \
          } else { \
-              RED_GFX7_PART(T,OP,ID); \
+             r = red7_part_##T##_##OP(x); \
          } \
     } else if (__oclc_ISA_version < 10000) { \
-        RED_GFX89(T,OP,ID); \
+        r = red89_##T##_##OP(x); \
     } else { \
-        RED_GFX10(T,OP,ID); \
+        r = red10_##T##_##OP(x); \
     } \
     return r; \
 }
 
-#define GENSCAN(T,OP,ID) \
+#define GENSCAN(T,OP,ID,IDZ) \
+GENISCAN7(T,OP,ID,IDZ) \
+GENISCAN89(T,OP,ID,IDZ) \
+GENISCAN10(T,OP,ID,IDZ) \
+GENSR1_7(T,OP,ID,IDZ) \
+GENSR1_89(T,OP,ID,IDZ) \
+GENSR1_10(T,OP,ID,IDZ) \
 IATTR T \
 C(__ockl_wfscan_,C(OP,T##_suf))(T x, bool inclusive) \
 { \
@@ -500,66 +551,66 @@ C(__ockl_wfscan_,C(OP,T##_suf))(T x, bool inclusive) \
     uint l = __ockl_lane_u32(); \
  \
     if (__oclc_ISA_version < 8000) { \
-        ISCAN_GFX7(T,OP,ID); \
+        s = iscan7_##T##_##OP(x, l); \
     } else  if (__oclc_ISA_version < 10000)  { \
-        ISCAN_GFX89(T,OP,ID); \
+        s = iscan89_##T##_##OP(x, l); \
     } else { \
-        ISCAN_GFX10(T,OP,ID); \
+        s = iscan10_##T##_##OP(x, l); \
     } \
  \
     if (!inclusive) { \
         if (__oclc_ISA_version < 8000) { \
-            SR1_SWIZZLE(T,ID); \
+            s = sr1_7_##T##_##OP(s, l); \
         } else  if (__oclc_ISA_version < 10000)  { \
-            SR1_GFX89(T,ID); \
+            s = sr1_89_##T##_##OP(s, l); \
         } else { \
-            SR1_GFX10(T,ID); \
+            s = sr1_10_##T##_##OP(s, l); \
         } \
     } \
  \
     return s; \
 }
 
-#define GEN(T,OP,ID) \
-    GENRED(T,OP,ID) \
-    GENSCAN(T,OP,ID)
-
-GEN(int,add,0)
-GEN(uint,add,0u)
-GEN(long,add,0L)
-GEN(ulong,add,0UL)
-GEN(float,add,0.0f)
-GEN(double,add,0.0)
-GEN(half,add,0.0h)
-
-GEN(int,min,INT_MAX)
-GEN(uint,min,UINT_MAX)
-GEN(long,min,LONG_MAX)
-GEN(ulong,min,ULONG_MAX)
-GEN(float,min,INFINITY)
-GEN(double,min,(double)INFINITY)
-GEN(half,min,(half)INFINITY)
-
-GEN(int,max,INT_MIN)
-GEN(uint,max,0u)
-GEN(long,max,LONG_MIN)
-GEN(ulong,max,0UL)
-GEN(float,max,-INFINITY)
-GEN(double,max,-(double)INFINITY)
-GEN(half,max,-(half)INFINITY)
-
-GEN(int,and,~0)
-GEN(uint,and,~0u)
-GEN(long,and,~0L)
-GEN(ulong,and,~0UL)
-
-GEN(int,or,0)
-GEN(uint,or,0u)
-GEN(long,or,0L)
-GEN(ulong,or,0UL)
-
-GEN(int,xor,0)
-GEN(uint,xor,0u)
-GEN(long,xor,0L)
-GEN(ulong,xor,0UL)
+#define GEN(T,OP,ID,IDZ) \
+    GENRED(T,OP,ID,IDZ) \
+    GENSCAN(T,OP,ID,IDZ)
+
+GEN(int,add,0,1)
+GEN(uint,add,0u,1)
+GEN(long,add,0L,1)
+GEN(ulong,add,0UL,1)
+GEN(float,add,0.0f,1)
+GEN(double,add,0.0,1)
+GEN(half,add,0.0h,1)
+
+GEN(int,min,INT_MAX,0)
+GEN(uint,min,UINT_MAX,0)
+GEN(long,min,LONG_MAX,0)
+GEN(ulong,min,ULONG_MAX,0)
+GEN(float,min,INFINITY,0)
+GEN(double,min,(double)INFINITY,0)
+GEN(half,min,(half)INFINITY,0)
+
+GEN(int,max,INT_MIN,0)
+GEN(uint,max,0u,1)
+GEN(long,max,LONG_MIN,0)
+GEN(ulong,max,0UL,1)
+GEN(float,max,-INFINITY,0)
+GEN(double,max,-(double)INFINITY,0)
+GEN(half,max,-(half)INFINITY,0)
+
+GEN(int,and,~0,0)
+GEN(uint,and,~0u,0)
+GEN(long,and,~0L,0)
+GEN(ulong,and,~0UL,0)
+
+GEN(int,or,0,1)
+GEN(uint,or,0u,1)
+GEN(long,or,0L,1)
+GEN(ulong,or,0UL,1)
+
+GEN(int,xor,0,1)
+GEN(uint,xor,0u,1)
+GEN(long,xor,0L,1)
+GEN(ulong,xor,0UL,1)