[SYCL][NFC] Fix SPIR-V format/spec name spelling (#2356)

SPIRV -> SPIR-V

Author: bader

sycl/doc/CompilerAndRuntimeDesign.md

@@ -21,7 +21,7 @@ target, then it will invoke the host compiler to compile the host part of a
 SYCL source. In the simplest case, when compilation and linkage are done in one
 compiler driver invocation, once compilation is finished, the device object
 files (which are really LLVM IR files) are linked with the `llvm-link` tool.
-The resulting LLVM IR module is then translated into a SPIRV module using the
+The resulting LLVM IR module is then translated into a SPIR-V module using the
 `llvm-spirv` tool and wrapped in a host object file using the
 `clang-offload-wrapper` tool. Once all the host object files and the wrapped
 object with device code are ready, the driver invokes the usual platform linker
@@ -32,7 +32,7 @@ line.
 There are many variations of the compilation process depending on whether user
 chose to do one or more of the following:
 - perform compilation separately from linkage
-- compile the device SPIRV module ahead-of-time for one or more targets
+- compile the device SPIR-V module ahead-of-time for one or more targets
 - perform device code splitting so that device code is distributed across
   multiple modules rather than enclosed in a single one
 - perform linkage of static device libraries
@@ -450,7 +450,7 @@ either from `llvm-spirv` or from the AOT backend.
 
 ##### Device code splitting
 
-Putting all device code into a single SPIRV module does not work well in the
+Putting all device code into a single SPIR-V module does not work well in the
 following cases:
 1. There are thousands of kernels defined and only small part of them is used at
 run-time. Having them all in one SPIR-V module significantly increases JIT time.

sycl/doc/extensions/ExplicitSIMD/dpcpp-explicit-simd.md

@@ -434,7 +434,7 @@ Explicit SIMD extension supports "private global" variables - file scope
 variables in private address space (similar to thread-local variables on host).
 These variables have 1 copy per work-item (which maps to a single SIMD thread in
 ESP) and are visible to all functions in the translation unit. Conceptually they
-map to SPIRV variable with private storage class. Private globals can be bound
+map to SPIR-V variable with private storage class. Private globals can be bound
 to a specific byte offset within the GRF. To mark a file scope variable as
 private global, the `INTEL_GPU_PRIVATE` attribute is used,
 `INTEL_GPU_REGISTER` is used to bind it the register file:

sycl/include/CL/sycl/INTEL/esimd/esimd_enum.hpp

@@ -26,8 +26,8 @@ using uint = unsigned int;
 // Mark a function being nodebug.
 #define ESIMD_NODEBUG __attribute__((nodebug))
 // Mark a "ESIMD global": accessible from all functions in current translation
-// unit, separate copy per subgroup (work-item), mapped to SPIRV private storage
-// class.
+// unit, separate copy per subgroup (work-item), mapped to SPIR-V private
+// storage class.
 #define ESIMD_PRIVATE __attribute__((opencl_private))
 // Bind a ESIMD global variable to a specific register.
 #define ESIMD_REGISTER(n) __attribute__((register_num(n)))

sycl/include/CL/sycl/ONEAPI/experimental/spec_constant.hpp

@@ -9,7 +9,7 @@
 // Based on:
 // https://github.com/codeplaysoftware/standards-proposals/blob/master/spec-constant/index.md
 // TODO:
-// 1) implement the SPIRV interop part of the proposal
+// 1) implement the SPIR-V interop part of the proposal
 // 2) support arbitrary spec constant type; only primitive types are
 //    supported currently
 // 3) move to the new upcoming spec constant specification (then 1 above is not

sycl/include/CL/sycl/detail/common.hpp

@@ -302,7 +302,7 @@ size_t getLinearIndex(const T<Dims> &Index, const U<Dims> &Range) {
 // pairs) into disjoint sets based on the kernel distribution among device
 // images.
 using KernelSetId = size_t;
-// Kernel set ID for kernels contained within the SPIRV file specified via
+// Kernel set ID for kernels contained within the SPIR-V file specified via
 // environment.
 constexpr KernelSetId SpvFileKSId = 0;
 constexpr KernelSetId LastKSId = SpvFileKSId;

sycl/include/CL/sycl/detail/image_ocl_types.hpp

@@ -95,15 +95,15 @@ static RetType __invoke__ImageReadSampler(ImageT Img, CoordT Coords,
 
   TempArgT TmpCoords =
       cl::sycl::detail::convertDataToType<CoordT, TempArgT>(Coords);
-  // According to validation rules(SPIRV specification, section 2.16.1) result
+  // According to validation rules(SPIR-V specification, section 2.16.1) result
   // of __spirv_SampledImage is allowed to be an operand of image lookup
   // and query instructions explicitly specified to take an operand whose
   // type is OpTypeSampledImage.
   //
-  // According to SPIRV specification section 3.32.10 at least one operand
+  // According to SPIR-V specification section 3.32.10 at least one operand
   // setting the level of detail must be present. The last two arguments of
   // __spirv_ImageSampleExplicitLod represent image operand type and value.
-  // From the SPIRV specification section 3.14:
+  // From the SPIR-V specification section 3.14:
   enum ImageOperands { Lod = 0x2 };
 
   // Lod value is zero as mipmap is not supported.
@@ -118,9 +118,9 @@ namespace sycl {
 namespace detail {
 
 // Function to return the number of channels for Image Channel Order returned by
-// SPIRV call to OpImageQueryOrder.
+// SPIR-V call to OpImageQueryOrder.
 // The returned int value represents an enum from Image Channel Order. The enums
-// for Image Channel Order are mapped differently in sycl and SPIRV spec.
+// for Image Channel Order are mapped differently in sycl and SPIR-V spec.
 inline int getSPIRVNumChannels(int ImageChannelOrder) {
   switch (ImageChannelOrder) {
   case 0:  // R
@@ -155,11 +155,11 @@ inline int getSPIRVNumChannels(int ImageChannelOrder) {
 }
 
 // Function to compute the Element Size for a given Image Channel Type and Image
-// Channel Order, returned by SPIRV calls to OpImageQueryFormat and
+// Channel Order, returned by SPIR-V calls to OpImageQueryFormat and
 // OpImageQueryOrder respectively.
 // The returned int value from OpImageQueryFormat represents an enum from Image
 // Channel Data Type. The enums for Image Channel Data Type are mapped
-// differently in sycl and SPIRV spec.
+// differently in sycl and SPIR-V spec.
 inline int getSPIRVElementSize(int ImageChannelType, int ImageChannelOrder) {
   int NumChannels = getSPIRVNumChannels(ImageChannelOrder);
   switch (ImageChannelType) {

sycl/include/CL/sycl/group.hpp

@@ -263,11 +263,11 @@ template <int Dimensions = 1> class group {
                      access::fence_space>::type accessSpace =
                      access::fence_space::global_and_local) const {
     uint32_t flags = detail::getSPIRVMemorySemanticsMask(accessSpace);
-    // TODO: currently, there is no good way in SPIRV to set the memory
+    // TODO: currently, there is no good way in SPIR-V to set the memory
     // barrier only for load operations or only for store operations.
     // The full read-and-write barrier is used and the template parameter
-    // 'accessMode' is ignored for now. Either SPIRV or SYCL spec may be
-    // changed to address this discrepancy between SPIRV and SYCL,
+    // 'accessMode' is ignored for now. Either SPIR-V or SYCL spec may be
+    // changed to address this discrepancy between SPIR-V and SYCL,
     // or if we decide that 'accessMode' is the important feature then
     // we can fix this later, for example, by using OpenCL 1.2 functions
     // read_mem_fence() and write_mem_fence().

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -814,7 +814,7 @@ pi_result piDeviceGetInfo(pi_device Device, pi_device_info ParamName,
     //
     std::string SupportedExtensions;
 
-    // cl_khr_il_program - OpenCL 2.0 KHR extension for SPIRV support. Core
+    // cl_khr_il_program - OpenCL 2.0 KHR extension for SPIR-V support. Core
     //   feature in >OpenCL 2.1
     // cl_khr_subgroups - Extension adds support for implementation-controlled
     //   subgroups.

sycl/source/detail/program_impl.cpp

@@ -506,7 +506,7 @@ void program_impl::flush_spec_constants(const RTDeviceBinaryImage &Img,
     const char *SCName = (*SCIt)->Name;
     auto SCEntry = SpecConstRegistry.find(SCName);
     if (SCEntry == SpecConstRegistry.end())
-      // spec constant has not been set in user code - SPIRV will use default
+      // spec constant has not been set in user code - SPIR-V will use default
       continue;
     const spec_constant_impl &SC = SCEntry->second;
     assert(SC.isSet() && "uninitialized spec constant");

sycl/source/detail/program_impl.hpp

@@ -298,7 +298,7 @@ class program_impl {
   /// managemment PI APIs. The native program passed as non-null argument
   /// overrides the MProgram native program field.
   /// \param Img device binary image corresponding to this program, used to
-  ///        resolve spec constant name to SPIRV integer ID
+  ///        resolve spec constant name to SPIR-V integer ID
   /// \param NativePrg if not null, used as the flush target, otherwise MProgram
   ///        is used
   void flush_spec_constants(const RTDeviceBinaryImage &Img,
@@ -417,7 +417,7 @@ class program_impl {
   OSModuleHandle MProgramModuleHandle = OSUtil::ExeModuleHandle;
 
   // Keeps specialization constant map for this program. Spec constant name
-  // resolution to actual SPIRV integer ID happens at build time, where the
+  // resolution to actual SPIR-V integer ID happens at build time, where the
   // device binary image is available. Access is guarded by this context's
   // program cache lock.
   SpecConstRegistryT SpecConstRegistry;

sycl/source/detail/program_manager/program_manager.cpp

@@ -240,7 +240,7 @@ static bool isDeviceBinaryTypeSupported(const context &C,
   const backend ContextBackend =
       detail::getSyclObjImpl(C)->getPlugin().getBackend();
 
-  // The CUDA backend cannot use SPIRV
+  // The CUDA backend cannot use SPIR-V
   if (ContextBackend == backend::cuda && Format == PI_DEVICE_BINARY_TYPE_SPIRV)
     return false;
 
@@ -379,7 +379,7 @@ RT::PiProgram ProgramManager::getBuiltPIProgram(OSModuleHandle M,
     // Link a fallback implementation of device libraries if they are not
     // supported by a device compiler.
     // Pre-compiled programs are supposed to be already linked.
-    // If device image is not SPIRV, DeviceLibReqMask will be 0 which means
+    // If device image is not SPIR-V, DeviceLibReqMask will be 0 which means
     // no fallback device library will be linked.
     uint32_t DeviceLibReqMask = 0;
     if (Img.getFormat() == PI_DEVICE_BINARY_TYPE_SPIRV &&
@@ -603,11 +603,12 @@ static RT::PiProgram loadDeviceLibFallback(
 
 ProgramManager::ProgramManager() {
   const char *SpvFile = std::getenv(UseSpvEnv);
-  // If a SPIRV file is specified with an environment variable,
+  // If a SPIR-V file is specified with an environment variable,
   // register the corresponding image
   if (SpvFile) {
     m_UseSpvFile = true;
-    // The env var requests that the program is loaded from a SPIRV file on disk
+    // The env var requests that the program is loaded from a SPIR-V file on
+    // disk
     std::ifstream File(SpvFile, std::ios::binary);
 
     if (!File.is_open())
@@ -805,7 +806,7 @@ ProgramManager::build(ProgramPtr Program, const ContextImplPtr Context,
 
   // TODO: this is a temporary workaround for GPU tests for ESIMD compiler.
   // We do not link with other device libraries, because it may fail
-  // due to unrecognized SPIRV format of those libraries.
+  // due to unrecognized SPIR-V format of those libraries.
   if (std::string(CompileOpts).find(std::string("-cmc")) != std::string::npos ||
       std::string(CompileOpts).find(std::string("-vc-codegen")) !=
           std::string::npos)

sycl/source/detail/program_manager/program_manager.hpp

@@ -192,7 +192,7 @@ class ProgramManager {
 
   // Keeps track of pi_program to image correspondence. Needed for:
   // - knowing which specialization constants are used in the program and
-  //   injecting their current values before compiling the SPIRV; the binary
+  //   injecting their current values before compiling the SPIR-V; the binary
   //   image object has info about all spec constants used in the module
   // - finding kernel argument masks for kernels associated with each
   //   pi_program
@@ -215,7 +215,7 @@ class ProgramManager {
   std::unordered_map<const RTDeviceBinaryImage *, KernelNameToArgMaskMap>
       m_EliminatedKernelArgMasks;
 
-  /// True iff a SPIRV file has been specified with an environment variable
+  /// True iff a SPIR-V file has been specified with an environment variable
   bool m_UseSpvFile = false;
 };
 } // namespace detail

sycl/source/detail/scheduler/commands.cpp

@@ -1628,7 +1628,7 @@ static void adjustNDRangePerKernel(NDRDescT &NDR, RT::PiKernel Kernel,
   NDR.set(NDR.Dims, nd_range<3>(NDR.NumWorkGroups * WGSize, WGSize));
 }
 
-// We have the following mapping between dimensions with SPIRV builtins:
+// We have the following mapping between dimensions with SPIR-V builtins:
 // 1D: id[0] -> x
 // 2D: id[0] -> y, id[1] -> x
 // 3D: id[0] -> z, id[1] -> y, id[2] -> x

sycl/source/detail/scheduler/commands.hpp

@@ -250,7 +250,7 @@ class Command {
   ///
   /// Stream ids are positive integers and we set it to an invalid value.
   int32_t MStreamID = -1;
-  /// Reserved for storing the object address such as SPIRV or memory object
+  /// Reserved for storing the object address such as SPIR-V or memory object
   /// address.
   void *MAddress = nullptr;
   /// Buffer to build the address string.

sycl/test/esimd/spirv_intrins_trans.cpp

@@ -1,5 +1,5 @@
 // RUN: %clangxx -fsycl -fsycl-explicit-simd -fsycl-device-only -O0 -S -emit-llvm -x c++ %s -o - | FileCheck %s
-// This test checks that all SPIRV intrinsics are correctly
+// This test checks that all SPIR-V intrinsics are correctly
 // translated into GenX counterparts (implemented in LowerCM.cpp)
 
 #include <CL/sycl.hpp>

sycl/test/spec_const/spec_const_types.cpp

@@ -12,7 +12,7 @@
 //===----------------------------------------------------------------------===//
 // The test checks that the tool chain correctly identifies all specialization
 // constants, emits correct specialization constats map file and can properly
-// translate the resulting bitcode to SPIRV.
+// translate the resulting bitcode to SPIR-V.
 
 #include <CL/sycl.hpp>