[SYCL][Docs] Add sycl_ext_oneapi_ternary_bitwise extension

clang/lib/Driver/ToolChains/Clang.cpp

@@ -10852,7 +10852,8 @@ static void getTripleBasedSPIRVTransOpts(Compilation &C,
       ",+SPV_INTEL_fpga_latency_control"
       ",+SPV_KHR_shader_clock"
       ",+SPV_INTEL_bindless_images"
-      ",+SPV_INTEL_task_sequence";
+      ",+SPV_INTEL_task_sequence"
+      ",+SPV_INTEL_ternary_bitwise_function";
   ExtArg = ExtArg + DefaultExtArg + INTELExtArg;
   if (C.getDriver().IsFPGAHWMode())
     // Enable several extensions on FPGA H/W exclusively

clang/lib/Sema/SPIRVBuiltins.td

@@ -901,6 +901,9 @@ foreach name = ["BitCount"] in {
   def : SPVBuiltin<name, [AIGenTypeN, AIGenTypeN], Attr.Const>;
 }
 
+def : SPVBuiltin<"BitwiseFunctionINTEL",
+                 [AIGenTypeN, AIGenTypeN, AIGenTypeN, AIGenTypeN, Int]>;
+
 // 3.32.20. Barrier Instructions
 
 foreach name = ["ControlBarrier"] in {

clang/test/Driver/sycl-spirv-ext-old-model.c

@@ -36,6 +36,7 @@
 // CHECK-DEFAULT-SAME:,+SPV_KHR_shader_clock
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_bindless_images
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_task_sequence
+// CHECK-DEFAULT-SAME:,+SPV_INTEL_ternary_bitwise_function
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_bfloat16_conversion
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_joint_matrix
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_hw_thread_queries
@@ -73,6 +74,7 @@
 // CHECK-CPU-SAME:,+SPV_INTEL_fpga_invocation_pipelining_attributes
 // CHECK-CPU-SAME:,+SPV_INTEL_fpga_latency_control
 // CHECK-CPU-SAME:,+SPV_INTEL_task_sequence
+// CHECK-CPU-SAME:,+SPV_INTEL_ternary_bitwise_function
 // CHECK-CPU-SAME:,+SPV_INTEL_bfloat16_conversion
 // CHECK-CPU-SAME:,+SPV_INTEL_joint_matrix
 // CHECK-CPU-SAME:,+SPV_INTEL_hw_thread_queries

clang/test/Driver/sycl-spirv-ext.c

@@ -53,6 +53,7 @@
 // CHECK-DEFAULT-SAME:,+SPV_KHR_shader_clock
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_bindless_images
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_task_sequence
+// CHECK-DEFAULT-SAME:,+SPV_INTEL_ternary_bitwise_function
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_bfloat16_conversion
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_joint_matrix
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_hw_thread_queries
@@ -90,6 +91,7 @@
 // CHECK-CPU-SAME:,+SPV_INTEL_fpga_invocation_pipelining_attributes
 // CHECK-CPU-SAME:,+SPV_INTEL_fpga_latency_control
 // CHECK-CPU-SAME:,+SPV_INTEL_task_sequence
+// CHECK-CPU-SAME:,+SPV_INTEL_ternary_bitwise_function
 // CHECK-CPU-SAME:,+SPV_INTEL_bfloat16_conversion
 // CHECK-CPU-SAME:,+SPV_INTEL_joint_matrix
 // CHECK-CPU-SAME:,+SPV_INTEL_hw_thread_queries

clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp

@@ -873,7 +873,8 @@ getTripleBasedSPIRVTransOpts(const ArgList &Args,
       ",+SPV_INTEL_fpga_latency_control"
       ",+SPV_KHR_shader_clock"
       ",+SPV_INTEL_bindless_images"
-      ",+SPV_INTEL_task_sequence";
+      ",+SPV_INTEL_task_sequence"
+      ",+SPV_INTEL_ternary_bitwise_function";
   ExtArg = ExtArg + DefaultExtArg + INTELExtArg;
   ExtArg += ",+SPV_INTEL_bfloat16_conversion"
             ",+SPV_INTEL_joint_matrix"

clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp

@@ -355,7 +355,8 @@ static void getSPIRVTransOpts(const ArgList &Args,
       ",+SPV_INTEL_fpga_latency_control"
       ",+SPV_INTEL_task_sequence"
       ",+SPV_KHR_shader_clock"
-      ",+SPV_INTEL_bindless_images";
+      ",+SPV_INTEL_bindless_images"
+      ",+SPV_INTEL_ternary_bitwise_function";
   ExtArg = ExtArg + DefaultExtArg + INTELExtArg;
   ExtArg += ",+SPV_INTEL_token_type"
             ",+SPV_INTEL_bfloat16_conversion"

llvm-spirv/include/LLVMSPIRVExtensions.inc

@@ -76,3 +76,4 @@ EXT(SPV_INTEL_maximum_registers)
 EXT(SPV_INTEL_bindless_images)
 EXT(SPV_INTEL_2d_block_io)
 EXT(SPV_INTEL_subgroup_matrix_multiply_accumulate)
+EXT(SPV_INTEL_ternary_bitwise_function)

llvm-spirv/lib/SPIRV/libSPIRV/SPIRVInstruction.h

@@ -4446,5 +4446,60 @@ class SPIRVSubgroupMatrixMultiplyAccumulateINTELInst
 _SPIRV_OP(SubgroupMatrixMultiplyAccumulate, true, 7, true, 4)
 #undef _SPIRV_OP
 
+class SPIRVTernaryBitwiseFunctionINTELInst : public SPIRVInstTemplateBase {
+public:
+  void validate() const override {
+    SPIRVInstruction::validate();
+    SPIRVErrorLog &SPVErrLog = this->getModule()->getErrorLog();
+    std::string InstName = "BitwiseFunctionINTEL";
+
+    const SPIRVType *ResTy = this->getType();
+    SPVErrLog.checkError(
+        ResTy->isTypeInt() || (ResTy->isTypeVector() &&
+                               ResTy->getVectorComponentType()->isTypeInt()),
+        SPIRVEC_InvalidInstruction,
+        InstName + "\nResult type must be an integer scalar or vector.\n");
+
+    auto CommonArgCheck = [this, ResTy, &InstName,
+                           &SPVErrLog](size_t ArgI, const char *ArgPlacement) {
+      SPIRVValue *Arg =
+          const_cast<SPIRVTernaryBitwiseFunctionINTELInst *>(this)->getOperand(
+              ArgI);
+      SPVErrLog.checkError(
+          Arg->getType() == ResTy, SPIRVEC_InvalidInstruction,
+          InstName + "\n" + ArgPlacement +
+              " argument must be the same as the result type.\n");
+    };
+
+    CommonArgCheck(0, "First");
+    CommonArgCheck(1, "Second");
+    CommonArgCheck(2, "Third");
+
+    SPIRVValue *LUTIndexArg =
+        const_cast<SPIRVTernaryBitwiseFunctionINTELInst *>(this)->getOperand(3);
+    const SPIRVType *LUTIndexArgTy = LUTIndexArg->getType();
+    SPVErrLog.checkError(
+        LUTIndexArgTy->isTypeInt(32), SPIRVEC_InvalidInstruction,
+        InstName + "\nFourth argument must be a 32-bit integer scalar.\n");
+    SPVErrLog.checkError(
+        isConstantOpCode(LUTIndexArg->getOpCode()), SPIRVEC_InvalidInstruction,
+        InstName + "\nFourth argument must be constant instruction.\n");
+  }
+
+  std::optional<ExtensionID> getRequiredExtension() const override {
+    return ExtensionID::SPV_INTEL_ternary_bitwise_function;
+  }
+  SPIRVCapVec getRequiredCapability() const override {
+    return getVec(internal::CapabilityTernaryBitwiseFunctionINTEL);
+  }
+};
+
+#define _SPIRV_OP(x, ...)                                                      \
+  typedef SPIRVInstTemplate<SPIRVTernaryBitwiseFunctionINTELInst,              \
+                            internal::Op##x##INTEL, __VA_ARGS__>               \
+      SPIRV##x##INTEL;
+_SPIRV_OP(BitwiseFunction, true, 7)
+#undef _SPIRV_OP
+
 } // namespace SPIRV
 #endif // SPIRV_LIBSPIRV_SPIRVINSTRUCTION_H

llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h

@@ -696,6 +696,8 @@ template <> inline void SPIRVMap<Capability, std::string>::init() {
       "SubgroupRequirementsINTEL");
   add(internal::CapabilityTaskSequenceINTEL, "TaskSequenceINTEL");
   add(internal::CapabilityBindlessImagesINTEL, "BindlessImagesINTEL");
+  add(internal::CapabilityTernaryBitwiseFunctionINTEL,
+      "TernaryBitwiseFunctionINTEL");
 }
 SPIRV_DEF_NAMEMAP(Capability, SPIRVCapabilityNameMap)
 

llvm-spirv/lib/SPIRV/libSPIRV/SPIRVOpCodeEnumInternal.h

@@ -44,3 +44,4 @@ _SPIRV_OP_INTERNAL(ConvertHandleToSamplerINTEL,
                    internal::ConvertHandleToSamplerINTEL)
 _SPIRV_OP_INTERNAL(ConvertHandleToSampledImageINTEL,
                    internal::ConvertHandleToSampledImageINTEL)
+_SPIRV_OP_INTERNAL(BitwiseFunctionINTEL, internal::BitwiseFunctionINTEL)

llvm-spirv/lib/SPIRV/libSPIRV/spirv_internal.hpp

@@ -89,6 +89,7 @@ enum InternalOp {
   IOpConvertHandleToImageINTEL = 6529,
   IOpConvertHandleToSamplerINTEL = 6530,
   IOpConvertHandleToSampledImageINTEL = 6531,
+  IOpBitwiseFunctionINTEL = 6242,
   IOpPrev = OpMax - 2,
   IOpForward
 };
@@ -124,7 +125,8 @@ enum InternalCapability {
   ICapabilityJointMatrixPackedInt4ComponentTypeINTEL = 6439,
   ICapabilityCacheControlsINTEL = 6441,
   ICapabilitySubgroupRequirementsINTEL = 6445,
-  ICapabilityBindlessImagesINTEL = 6528
+  ICapabilityBindlessImagesINTEL = 6528,
+  ICapabilityTernaryBitwiseFunctionINTEL = 6241
 };
 
 enum InternalFunctionControlMask { IFunctionControlOptNoneINTELMask = 0x10000 };
@@ -222,6 +224,9 @@ _SPIRV_OP(Capability, BindlessImagesINTEL)
 _SPIRV_OP(Op, ConvertHandleToImageINTEL)
 _SPIRV_OP(Op, ConvertHandleToSamplerINTEL)
 _SPIRV_OP(Op, ConvertHandleToSampledImageINTEL)
+
+_SPIRV_OP(Capability, TernaryBitwiseFunctionINTEL)
+_SPIRV_OP(Op, BitwiseFunctionINTEL)
 #undef _SPIRV_OP
 
 constexpr SourceLanguage SourceLanguagePython =

llvm-spirv/test/extensions/INTEL/SPV_INTEL_ternary_bitwise_function/bitwise_function.ll

@@ -0,0 +1,78 @@
+; RUN: llvm-as %s -o %t.bc
+; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_ternary_bitwise_function -o %t.spv
+; RUN: llvm-spirv %t.spv --to-text -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
+
+; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
+; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
+
+; RUN: not llvm-spirv %t.bc 2>&1 | FileCheck %s --check-prefix=CHECK-ERROR
+; CHECK-ERROR: RequiresExtension: Feature requires the following SPIR-V extension:
+; CHECK-ERROR-NEXT: SPV_INTEL_ternary_bitwise_function
+
+; CHECK-SPIRV-NOT: Name [[#]] "_Z28__spirv_BitwiseFunctionINTELiiij"
+; CHECK-SPIRV-NOT: Name [[#]] "_Z28__spirv_BitwiseFunctionINTELDv4_iS_S_j"
+
+; CHECK-SPIRV-DAG: Capability TernaryBitwiseFunctionINTEL
+; CHECK-SPIRV-DAG: Extension "SPV_INTEL_ternary_bitwise_function"
+
+; CHECK-SPIRV-DAG: TypeInt [[#TYPEINT:]] 32 0
+; CHECK-SPIRV-DAG: TypeVector [[#TYPEINTVEC4:]] [[#TYPEINT]] 4
+; CHECK-SPIRV-DAG: Constant [[#TYPEINT]] [[#ScalarLUT:]] 24
+; CHECK-SPIRV-DAG: Constant [[#TYPEINT]] [[#VecLUT:]] 42
+
+; CHECK-SPIRV: Load [[#TYPEINT]] [[#ScalarA:]]
+; CHECK-SPIRV: Load [[#TYPEINT]] [[#ScalarB:]]
+; CHECK-SPIRV: Load [[#TYPEINT]] [[#ScalarC:]]
+; CHECK-SPIRV: BitwiseFunctionINTEL [[#TYPEINT]] {{.*}} [[#ScalarA]] [[#ScalarB]] [[#ScalarC]] [[#ScalarLUT]]
+; CHECK-SPIRV: Load [[#TYPEINTVEC4]] [[#VecA:]]
+; CHECK-SPIRV: Load [[#TYPEINTVEC4]] [[#VecB:]]
+; CHECK-SPIRV: Load [[#TYPEINTVEC4]] [[#VecC:]]
+; CHECK-SPIRV: BitwiseFunctionINTEL [[#TYPEINTVEC4]] {{.*}} [[#VecA]] [[#VecB]] [[#VecC]] [[#VecLUT]]
+
+; CHECK-LLVM: %[[ScalarA:.*]] = load i32, ptr
+; CHECK-LLVM: %[[ScalarB:.*]] = load i32, ptr
+; CHECK-LLVM: %[[ScalarC:.*]] = load i32, ptr
+; CHECK-LLVM: call spir_func i32 @_Z28__spirv_BitwiseFunctionINTELiiii(i32 %[[ScalarA]], i32 %[[ScalarB]], i32 %[[ScalarC]], i32 24)
+; CHECK-LLVM: %[[VecA:.*]] = load <4 x i32>, ptr
+; CHECK-LLVM: %[[VecB:.*]] = load <4 x i32>, ptr
+; CHECK-LLVM: %[[VecC:.*]] = load <4 x i32>, ptr
+; CHECK-LLVM: call spir_func <4 x i32> @_Z28__spirv_BitwiseFunctionINTELDv4_iS_S_i(<4 x i32> %[[VecA]], <4 x i32> %[[VecB]], <4 x i32> %[[VecC]], i32 42)
+
+target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
+target triple = "spir"
+
+; Function Attrs: nounwind readnone
+define spir_kernel void @fooScalar() {
+entry:
+  %argA = alloca i32
+  %argB = alloca i32
+  %argC = alloca i32
+  %A = load i32, ptr %argA
+  %B = load i32, ptr %argB
+  %C = load i32, ptr %argC
+  %res = call spir_func i32 @_Z28__spirv_BitwiseFunctionINTELiiii(i32 %A, i32 %B, i32 %C, i32 24)
+  ret void
+}
+
+; Function Attrs: nounwind readnone
+define spir_kernel void @fooVec() {
+entry:
+  %argA = alloca <4 x i32>
+  %argB = alloca <4 x i32>
+  %argC = alloca <4 x i32>
+  %A = load <4 x i32>, ptr %argA
+  %B = load <4 x i32>, ptr %argB
+  %C = load <4 x i32>, ptr %argC
+  %res = call spir_func <4 x i32> @_Z28__spirv_BitwiseFunctionINTELDv4_iS_S_i(<4 x i32> %A, <4 x i32> %B, <4 x i32> %C, i32 42)
+  ret void
+}
+
+declare dso_local spir_func i32 @_Z28__spirv_BitwiseFunctionINTELiiii(i32, i32, i32, i32)
+declare dso_local spir_func <4 x i32> @_Z28__spirv_BitwiseFunctionINTELDv4_iS_S_i(<4 x i32>, <4 x i32>, <4 x i32>, i32)
+
+!llvm.module.flags = !{!0}
+!opencl.spir.version = !{!1}
+
+!0 = !{i32 1, !"wchar_size", i32 4}
+!1 = !{i32 1, i32 2}

sycl/doc/extensions/experimental/sycl_ext_oneapi_ternary_bitwise.asciidoc

@@ -0,0 +1,115 @@
+= sycl_ext_oneapi_ternary_bitwise
+
+:source-highlighter: coderay
+:coderay-linenums-mode: table
+
+// This section needs to be after the document title.
+:doctype: book
+:toc2:
+:toc: left
+:encoding: utf-8
+:lang: en
+:dpcpp: pass:[DPC++]
+:endnote: —{nbsp}end{nbsp}note
+
+// Set the default source code type in this document to C++,
+// for syntax highlighting purposes.  This is needed because
+// docbook uses c++ and html5 uses cpp.
+:language: {basebackend@docbook:c++:cpp}
+
+
+== Notice
+
+[%hardbreaks]
+Copyright (C) 2025 Intel Corporation.  All rights reserved.
+
+Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks
+of The Khronos Group Inc.  OpenCL(TM) is a trademark of Apple Inc. used by
+permission by Khronos.
+
+
+== Contact
+
+To report problems with this extension, please open a new issue at:
+
+https://github.com/intel/llvm/issues
+
+
+== Dependencies
+
+This extension is written against the SYCL 2020 revision 9 specification.  All
+references below to the "core SYCL specification" or to section numbers in the
+SYCL specification refer to that revision.
+
+
+== Status
+
+This is an experimental extension specification, intended to provide early
+access to features and gather community feedback.  Interfaces defined in this
+specification are implemented in {dpcpp}, but they are not finalized and may
+change incompatibly in future versions of {dpcpp} without prior notice.
+*Shipping software products should not rely on APIs defined in this
+specification.*
+
+
+== Overview
+
+Some hardware offers efficient bitwise operations on three arguments. To expose
+these bitwise operations in SYCL, this extension adds a new `ternary_bitwise`
+function, where the bitwise operation computed is controlled through a
+look-up table (LUT) index computed by applying the bitwise operation to a
+predefined set of operands.
+
+
+== Specification
+
+=== Feature test macro
+
+This extension provides a feature-test macro as described in the core SYCL
+specification.  An implementation supporting this extension must predefine the
+macro `SYCL_EXT_ONEAPI_TERNARY_BITWISE` to one of the values defined in the table
+below.  Applications can test for the existence of this macro to determine if
+the implementation supports this feature, or applications can test the macro's
+value to determine which of the extension's features the implementation
+supports.
+
+[%header,cols="1,5"]
+|===
+|Value
+|Description
+
+|1
+|The APIs of this experimental extension are not versioned, so the
+ feature-test macro always has this value.
+|===
+
+=== New `ternary_bitwise` function
+
+
+|====
+a|
+[frame=all,grid=none]
+!====
+a!
+[source]
+----
+namespace sycl::ext::oneapi::experimental {
+
+template <uint8_t LUTIndex, typename T> T ternary_bitwise(T A, T B, T C)
+
+} // namespace sycl::ext::oneapi::experimental
+----
+!====
+
+_Constraints:_ The type `T` must be a generic integer type, as listed in section
+link:https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:integer-functions[4.17.7. Integer functions]
+of the SYCL 2020 specification.
+
+_Returns:_ The result of the bitwise operation identified by `LUTIndex` applied
+to `a`, `b` and `c`.  For any `I` it holds that
+`ternary_bitwise<I, uint8_t>(0xAA, 0xCC, 0xF0) == I` evaluates to `true, so the
+`LUTIndex` for a bitwise operation can be determined by applying the bitwise
+operation to `0xAA`, `0xCC` and `0xF0` in place of `a`, `b` and `c`
+respectively.
+|====
+