diff --git a/sycl/doc/CompilerAndRuntimeDesign.md b/sycl/doc/CompilerAndRuntimeDesign.md index 4fae3b6fb4880..5eb837f4159ba 100644 --- a/sycl/doc/CompilerAndRuntimeDesign.md +++ b/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. diff --git a/sycl/doc/extensions/ExplicitSIMD/dpcpp-explicit-simd.md b/sycl/doc/extensions/ExplicitSIMD/dpcpp-explicit-simd.md index 501d19b172e16..6093db83d9a93 100644 --- a/sycl/doc/extensions/ExplicitSIMD/dpcpp-explicit-simd.md +++ b/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: diff --git a/sycl/include/CL/sycl/INTEL/esimd/esimd_enum.hpp b/sycl/include/CL/sycl/INTEL/esimd/esimd_enum.hpp index 78ceb5c0c4b97..261b34391b9c7 100644 --- a/sycl/include/CL/sycl/INTEL/esimd/esimd_enum.hpp +++ b/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))) diff --git a/sycl/include/CL/sycl/ONEAPI/experimental/spec_constant.hpp b/sycl/include/CL/sycl/ONEAPI/experimental/spec_constant.hpp index 03698008e056d..bd8f4f9a3fb57 100644 --- a/sycl/include/CL/sycl/ONEAPI/experimental/spec_constant.hpp +++ b/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 diff --git a/sycl/include/CL/sycl/detail/common.hpp b/sycl/include/CL/sycl/detail/common.hpp index 11bb8f395a7c8..88ac3a4af8263 100644 --- a/sycl/include/CL/sycl/detail/common.hpp +++ b/sycl/include/CL/sycl/detail/common.hpp @@ -302,7 +302,7 @@ size_t getLinearIndex(const T &Index, const U &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; diff --git a/sycl/include/CL/sycl/detail/image_ocl_types.hpp b/sycl/include/CL/sycl/detail/image_ocl_types.hpp index 3055222755b5e..d6aa4e6e84963 100644 --- a/sycl/include/CL/sycl/detail/image_ocl_types.hpp +++ b/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(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) { diff --git a/sycl/include/CL/sycl/group.hpp b/sycl/include/CL/sycl/group.hpp index 7d7a8c504be46..9d474cd6187c1 100644 --- a/sycl/include/CL/sycl/group.hpp +++ b/sycl/include/CL/sycl/group.hpp @@ -263,11 +263,11 @@ template 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(). diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 2c6c95d2f6089..b7524128e50df 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/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. diff --git a/sycl/source/detail/program_impl.cpp b/sycl/source/detail/program_impl.cpp index 8453584aa067c..3e722494864cb 100644 --- a/sycl/source/detail/program_impl.cpp +++ b/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"); diff --git a/sycl/source/detail/program_impl.hpp b/sycl/source/detail/program_impl.hpp index b4153eb151339..c858004b46971 100644 --- a/sycl/source/detail/program_impl.hpp +++ b/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; diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index a4f2f875d424e..e3b9f01b7a3af 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/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) diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp index 1e4329b3fd925..d85d8ac3961a6 100644 --- a/sycl/source/detail/program_manager/program_manager.hpp +++ b/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 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 diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index efbe0373f06e6..f71a4e20df170 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/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 diff --git a/sycl/source/detail/scheduler/commands.hpp b/sycl/source/detail/scheduler/commands.hpp index 0936de076a6cc..1b551bce0c470 100644 --- a/sycl/source/detail/scheduler/commands.hpp +++ b/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. diff --git a/sycl/test/esimd/spirv_intrins_trans.cpp b/sycl/test/esimd/spirv_intrins_trans.cpp index 0f8eee3972a35..05745e1019c06 100644 --- a/sycl/test/esimd/spirv_intrins_trans.cpp +++ b/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 diff --git a/sycl/test/spec_const/spec_const_types.cpp b/sycl/test/spec_const/spec_const_types.cpp index 0a5178c5824eb..f5e4af764d7bb 100644 --- a/sycl/test/spec_const/spec_const_types.cpp +++ b/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