Optimize SPIR-V / OpenCL C "bitselect" builtin function

Add a dedicated intrinsic that guarantees "bitselect" builtin is
implemented with one "bfn" instruction.

Author: pkwasnie-intel

IGC/BiFModule/Headers/bif_flag_controls.h

@@ -23,6 +23,7 @@ BIF_FLAG_CONTROL(bool, UseNativeFP16AtomicMinMax)
 BIF_FLAG_CONTROL(bool, HasInt64SLMAtomicCAS)
 BIF_FLAG_CONTROL(bool, UseNativeFP64GlobalAtomicAdd)
 BIF_FLAG_CONTROL(bool, UseNative64BitIntBuiltin)
+BIF_FLAG_CONTROL(bool, UseBfn)
 BIF_FLAG_CONTROL(bool, HasThreadPauseSupport)
 BIF_FLAG_CONTROL(bool, UseNative64BitFloatBuiltin)
 BIF_FLAG_CONTROL(bool, hasHWLocalThreadID)

IGC/BiFModule/Implementation/IGCBiF_Intrinsics.cl

@@ -224,6 +224,10 @@ double    __builtin_IB_dmin(double, double) __attribute__((const));
 double    __builtin_IB_dmax(double, double) __attribute__((const));
 #endif
 
+// Boolean function on three sources
+short __builtin_IB_bfn_i16(short, short, short, uchar) __attribute__((const));
+int   __builtin_IB_bfn_i32(int, int, int, uchar) __attribute__((const));
+
 // Atomic operations
 int      __builtin_IB_atomic_add_global_i32(__global int*, int);
 int      __builtin_IB_atomic_add_local_i32(__local int*, int);

IGC/BiFModule/Implementation/Relational/bitselect.cl

@@ -9,43 +9,81 @@ SPDX-License-Identifier: MIT
 #include "../include/BiF_Definitions.cl"
 #include "../../Headers/spirv.h"
 
+// Bitselect can be implemented with the following boolean function:
+//   s0 & s1 | ~s0 & s2
+// where s0 = c, s1 = b, s2 = a
+// This maps to boolean function 0xD8.
 
 INLINE
 char SPIRV_OVERLOADABLE SPIRV_OCL_BUILTIN(bitselect, _i8_i8_i8, )( char a, char b, char c )
 {
-    char temp;
-    temp = (c & b) | (~c & a);
-    return temp;
+    if (BIF_FLAG_CTRL_GET(UseBfn))
+    {
+        return (char) __builtin_IB_bfn_i16((short)as_uchar(c), (short)as_uchar(b), (short)as_uchar(a), 0xD8);
+    }
+    else
+    {
+        char temp;
+        temp = (c & b) | (~c & a);
+        return temp;
+    }
 }
 
 GENERATE_SPIRV_OCL_VECTOR_FUNCTIONS_3ARGS( bitselect, char, char, i8 )
 
 INLINE
 short SPIRV_OVERLOADABLE SPIRV_OCL_BUILTIN(bitselect, _i16_i16_i16, )( short a, short b, short c )
 {
-    short temp;
-    temp = (c & b) | (~c & a);
-    return temp;
+    if (BIF_FLAG_CTRL_GET(UseBfn))
+    {
+        return __builtin_IB_bfn_i16(c, b, a, 0xD8);
+    }
+    else
+    {
+        short temp;
+        temp = (c & b) | (~c & a);
+        return temp;
+    }
 }
 
 GENERATE_SPIRV_OCL_VECTOR_FUNCTIONS_3ARGS( bitselect, short, short, i16 )
 
 INLINE
 int SPIRV_OVERLOADABLE SPIRV_OCL_BUILTIN(bitselect, _i32_i32_i32, )( int a, int b, int c )
 {
-    int temp;
-    temp = (c & b) | (~c & a);
-    return temp;
+    if (BIF_FLAG_CTRL_GET(UseBfn))
+    {
+        return __builtin_IB_bfn_i32(c, b, a, 0xD8);
+    }
+    else
+    {
+        int temp;
+        temp = (c & b) | (~c & a);
+        return temp;
+    }
 }
 
 GENERATE_SPIRV_OCL_VECTOR_FUNCTIONS_3ARGS( bitselect, int, int, i32 )
 
 INLINE
 long SPIRV_OVERLOADABLE SPIRV_OCL_BUILTIN(bitselect, _i64_i64_i64, )( long a, long b, long c )
 {
-    long temp;
-    temp = (c & b) | (~c & a);
-    return temp;
+    if (BIF_FLAG_CTRL_GET(UseBfn))
+    {
+        int2 tmpA = as_int2(a);
+        int2 tmpB = as_int2(b);
+        int2 tmpC = as_int2(c);
+        int2 tmpResult;
+        tmpResult.s0 = __builtin_IB_bfn_i32(tmpC.s0, tmpB.s0, tmpA.s0, 0xD8);
+        tmpResult.s1 = __builtin_IB_bfn_i32(tmpC.s1, tmpB.s1, tmpA.s1, 0xD8);
+        return as_long(tmpResult);
+    }
+    else
+    {
+        long temp;
+        temp = (c & b) | (~c & a);
+        return temp;
+    }
 }
 
 GENERATE_SPIRV_OCL_VECTOR_FUNCTIONS_3ARGS( bitselect, long, long, i64 )

IGC/Compiler/Builtins/BIFFlagCtrl/BIFFlagCtrlResolution.cpp

@@ -66,6 +66,8 @@ void BIFFlagCtrlResolution::FillFlagCtrl() {
                     PtrCGC->platform.hasThreadPauseSupport());
   BIF_FLAG_CTRL_SET(UseNative64BitFloatBuiltin,
                     !PtrCGC->platform.hasNoFP64Inst());
+  BIF_FLAG_CTRL_SET(UseBfn, IGC_IS_FLAG_ENABLED(EnableBfn) &&
+      PtrCGC->platform.supportBfnInstruction());
   BIF_FLAG_CTRL_SET(hasHWLocalThreadID, PtrCGC->platform.hasHWLocalThreadID());
   BIF_FLAG_CTRL_SET(CRMacros, PtrCGC->platform.hasCorrectlyRoundedMacros());
   BIF_FLAG_CTRL_SET(

IGC/Compiler/CISACodeGen/EmitVISAPass.cpp

@@ -9604,6 +9604,9 @@ void EmitPass::EmitGenIntrinsicMessage(llvm::GenIntrinsicInst* inst)
     case GenISAIntrinsic::GenISA_WaveClusteredInterleave:
         emitWaveClusteredInterleave(inst);
         break;
+    case GenISAIntrinsic::GenISA_bfn:
+        emitBfn(inst);
+        break;
     case GenISAIntrinsic::GenISA_dp4a_ss:
     case GenISAIntrinsic::GenISA_dp4a_uu:
     case GenISAIntrinsic::GenISA_dp4a_su:
@@ -17789,6 +17792,19 @@ void EmitPass::emitFPOWithNonDefaultRoundingMode(llvm::GenIntrinsicInst* inst)
     ResetRoundingMode(inst);
 }
 
+void EmitPass::emitBfn(llvm::GenIntrinsicInst* inst)
+{
+    IGC_ASSERT_MESSAGE(isa<ConstantInt>(inst->getArgOperand(3)), "booleanFuncCtrl must be const!");
+    const uint8_t booleanFuncCtrl = int_cast<uint8_t>(cast<ConstantInt>(inst->getArgOperand(3))->getZExtValue());
+
+    CVariable* src0 = GetSymbol(inst->getOperand(0));
+    CVariable* src1 = GetSymbol(inst->getOperand(1));
+    CVariable* src2 = GetSymbol(inst->getOperand(2));
+
+    m_encoder->Bfn(booleanFuncCtrl, m_destination, src0, src1, src2);
+    m_encoder->Push();
+}
+
 void EmitPass::emitftoi(llvm::GenIntrinsicInst* inst)
 {
     IGC_ASSERT_MESSAGE(inst->getOperand(0)->getType()->isFloatingPointTy(), "Unsupported type");

IGC/Compiler/CISACodeGen/EmitVISAPass.hpp

@@ -455,6 +455,8 @@ class EmitPass : public llvm::FunctionPass
     void emitftoi(llvm::GenIntrinsicInst* inst);
     void emitCtlz(const SSource& source);
 
+    void emitBfn(llvm::GenIntrinsicInst* inst);
+
 
     // VME
     void emitVMESendIME(llvm::GenIntrinsicInst* inst);

IGC/Compiler/CISACodeGen/WIAnalysis.cpp

@@ -1508,7 +1508,8 @@ WIAnalysis::WIDependancy WIAnalysisRunner::calculate_dep(const CallInst* inst)
         GII_id == GenISAIntrinsic::GenISA_LSC2DBlockCreateAddrPayload ||
         GII_id == GenISAIntrinsic::GenISA_LSC2DBlockCopyAddrPayload ||
         GII_id == GenISAIntrinsic::GenISA_PredicatedLoad ||
-        GII_id == GenISAIntrinsic::GenISA_PredicatedStore)
+        GII_id == GenISAIntrinsic::GenISA_PredicatedStore ||
+        GII_id == GenISAIntrinsic::GenISA_bfn)
     {
         switch (GII_id)
         {

IGC/Compiler/CISACodeGen/opCode.h

@@ -156,6 +156,7 @@ DECLARE_OPCODE(GenISA_dp4a_ss, GenISAIntrinsic, llvm_dp4a_ss, false, true, true,
 DECLARE_OPCODE(GenISA_dp4a_uu, GenISAIntrinsic, llvm_dp4a_uu, false, true, true, true, false, false, false)
 DECLARE_OPCODE(GenISA_dp4a_su, GenISAIntrinsic, llvm_dp4a_su, false, true, true, true, false, false, false)
 DECLARE_OPCODE(GenISA_dp4a_us, GenISAIntrinsic, llvm_dp4a_us, false, true, true, true, false, false, false)
+DECLARE_OPCODE(GenISA_bfn, GenISAIntrinsic, llvm_bfn, false, false, false, false, false, false, false)
 
 // GS Intrinsics
 DECLARE_OPCODE(GenISA_OUTPUTGS, GenISAIntrinsic, llvm_output_gs, false, false, false, false, false, false, false)

IGC/Compiler/Optimizer/OCLBIUtils.cpp

@@ -1758,6 +1758,9 @@ CBuiltinsResolver::CBuiltinsResolver(CImagesBI::ParamMap* paramMap, CImagesBI::I
 
     m_CommandMap[StringRef("__builtin_IB_samplepos")] = CSamplePos::create();
 
+    m_CommandMap["__builtin_IB_bfn_i16"] = CSimpleIntrinMapping::create(GenISAIntrinsic::GenISA_bfn);
+    m_CommandMap["__builtin_IB_bfn_i32"] = CSimpleIntrinMapping::create(GenISAIntrinsic::GenISA_bfn);
+
     // `dp4a` built-ins
     m_CommandMap["__builtin_IB_dp4a_ss"] = CSimpleIntrinMapping::create(GenISAIntrinsic::GenISA_dp4a_ss, false);
     m_CommandMap["__builtin_IB_dp4a_uu"] = CSimpleIntrinMapping::create(GenISAIntrinsic::GenISA_dp4a_uu, false);

IGC/GenISAIntrinsics/generator/input/Intrinsic_definitions.yml

@@ -11730,3 +11730,31 @@ intrinsics:
         memory_effects:
             - !<MemoryRestriction>
                 memory_access: !MemoryAccessType NoModRef
+    - !<IntrinsicDefinition>
+        name: "GenISA_bfn"
+        comment: "Performs an arbitrary boolean logical operation on three sources."
+        return_definition: !<ReturnDefinition>
+            type_definition: *any_int
+            comment: "result"
+        arguments:
+            - !<ArgumentDefinition>
+                name: Arg0
+                type_definition: *ref_0_
+                comment: "source0 (a)"
+            - !<ArgumentDefinition>
+                name: Arg1
+                type_definition: *ref_0_
+                comment: "source1 (b)"
+            - !<ArgumentDefinition>
+                name: Arg2
+                type_definition: *ref_0_
+                comment: "source2 (c)"
+            - !<ArgumentDefinition>
+                name: Arg3
+                type_definition: *i8
+                comment: "Boolean function"
+        attributes:
+            - !AttributeID "NoUnwind"
+        memory_effects:
+            - !<MemoryRestriction>
+                memory_access: !MemoryAccessType NoModRef

IGC/VectorCompiler/lib/GenXOpts/CMTrans/GenXBIFFlagCtrlResolution.cpp

@@ -89,6 +89,7 @@ void GenXBIFFlagCtrlResolution::FillFlagCtrl() {
   BIF_FLAG_CTRL_SET(hasHWLocalThreadID, false);
   BIF_FLAG_CTRL_SET(APIRS, false);
   BIF_FLAG_CTRL_SET(UseLSC, false);
+  BIF_FLAG_CTRL_SET(UseBfn, false);
   BIF_FLAG_CTRL_SET(ForceL1Prefetch, false);
   BIF_FLAG_CTRL_SET(UseNativeFP64GlobalAtomicAdd, false);
   BIF_FLAG_CTRL_SET(MaxHWThreadIDPerSubDevice, 1);

IGC/ocloc_tests/Builtins/bitselect.cl

@@ -0,0 +1,37 @@
+/*========================== begin_copyright_notice ============================
+
+Copyright (C) 2025 Intel Corporation
+
+SPDX-License-Identifier: MIT
+
+============================= end_copyright_notice ===========================*/
+
+// REQUIRES: regkeys, xe2-supported
+// UNSUPPORTED: sys32
+
+// RUN: ocloc compile -file %s -options " -igc_opts 'DumpVISAASMToConsole=1'" -device bmg | FileCheck %s
+
+// CHECK: .kernel "test_i32"
+// CHECK: lsc_load.ugm (M1, 32)  [[A:V[0-9]+]]:d32  flat[{{V[0-9]+}}]:a64
+// CHECK: lsc_load.ugm (M1, 32)  [[B:V[0-9]+]]:d32  flat[{{V[0-9]+}}]:a64
+// CHECK: lsc_load.ugm (M1, 32)  [[C:V[0-9]+]]:d32  flat[{{V[0-9]+}}]:a64
+// CHECK: bfn.xd8 (M1, 32) [[RESULT:V[0-9]+]](0,0)<1> [[C]](0,0)<1;1,0> [[B]](0,0)<1;1,0> [[A]](0,0)<1;1,0>
+// CHECK: lsc_store.ugm (M1, 32)  flat[{{V[0-9]+}}]:a64  [[RESULT]]:d32
+__attribute__((intel_reqd_sub_group_size(32)))
+kernel void test_i32(global int* a, global int* b, global int* c) {
+    int i = get_global_id(0);
+    a[i] = bitselect(a[i], b[i], c[i]);
+}
+
+// CHECK: .kernel "test_i64"
+// CHECK: lsc_load.ugm (M1, 32)  [[A:V[0-9]+]]:d32x2  flat[{{V[0-9]+}}]:a64
+// CHECK: lsc_load.ugm (M1, 32)  [[B:V[0-9]+]]:d32x2  flat[{{V[0-9]+}}]:a64
+// CHECK: lsc_load.ugm (M1, 32)  [[C:V[0-9]+]]:d32x2  flat[{{V[0-9]+}}]:a64
+// CHECK: bfn.xd8 (M1, 32) [[RESULT:.+]](0,0)<1> [[C]](0,0)<1;1,0> [[B]](0,0)<1;1,0> [[A]](0,0)<1;1,0>
+// CHECK: bfn.xd8 (M1, 32) [[RESULT]](2,0)<1> [[C]](2,0)<1;1,0> [[B]](2,0)<1;1,0> [[A]](2,0)<1;1,0>
+// CHECK: lsc_store.ugm (M1, 32)  flat[{{V[0-9]+}}]:a64  [[RESULT]]:d32x2
+__attribute__((intel_reqd_sub_group_size(32)))
+kernel void test_i64(global long* a, global long* b, global long* c) {
+    int i = get_global_id(0);
+    a[i] = bitselect(a[i], b[i], c[i]);
+}