SPIR-V*：面向 OpenCL™ 工作负载的英特尔® 显卡编译器默认接口-演道网

英特尔® 显卡编译器最近从 SPIR* 转换到 SPIR-V*，作为面向 OpenCL™ 工作负载的中间表示。这看起来像编译器的内部变化，对用户来说不可见，但是这展示了我们支持 Khronos* 开放标准的承诺。现在，下载至编译器的大多数计算工作负载都将使用 SPIR-V。

下文简要介绍了在 OpenCL 应用中使用 SPIR-V 的优势以及如何开始使用它。

SPIR-V* – 面向显卡和计算工作负载的 Khronos* 中间表示

中间表示 (IR) 在现代编译器架构中扮演重要角色。它们使人类可读的源语言（如 C++、GLSL 或 OpenCL* C）更接近机器表示，但是它们仍比较抽象。IR 支持编译器独立于目标机器规范执行常见的优化，或在英特尔® 显卡编译器中执行 GPU 生成。

英特尔® 显卡编译器是英特尔® 显卡驱动程序的一部分，旨在将各种前端（如 Vulkan*、OpenGL*、DirectX*）生成的此类中间表示用作输入，并输出面向特定 GPU 生成的机器代码。

SPIR-V 是一种 Khronos 支持的现代二进制低级中间表示语言，用于 Vulkan、OpenCL、OpenGL 等许多 API。¹在 OpenCL 世界中，它是基于 LLVM-IR 的 SPIR² 语言的后续产品。

如欲获取 SPIR-V 设计的概述，请查看 SPIR-V 白皮书。

由于 SPIR-V 是一种二进制格式，因此人类无法读取。尽管如此，我们来看一下使用 cl_intel_subgroups 扩展的简单的反汇编 OpenCL 内核:

OpenCL C 内核：

#pragma OPENCL EXTENSION cl_intel_subgroups : enable

__kernel void testKernel(__global uint *shuffle_results) {
    size_t tid = get_global_id(0);
    uint value = get_sub_group_local_id();
        uint new_value = intel_sub_group_shuffle(value, 1);
        shuffle_results[tid] = new_value;
}

生成的 SPIR-V：

; SPIR-V
; Version: 1.0
; Generator: Khronos LLVM/SPIR-V Translator; 14
; Bound: 36
; Schema: 0
               OpCapability Addresses
               OpCapability Linkage
               OpCapability Kernel
               OpCapability Int64
               OpCapability SubgroupShuffleINTEL
          %1 = OpExtInstImport "OpenCL.std"
               OpMemoryModel Physical64 OpenCL
               OpEntryPoint Kernel %12 "testKernel"
         %32 = OpString "kernel_arg_type.testKernel.uint*,"
               OpSource OpenCL_C 200000
               OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
               OpName %__spirv_BuiltInSubgroupLocalInvocationId "__spirv_BuiltInSubgroupLocalInvocationId"
               OpName %shuffle_results "shuffle_results"
               OpName %shuffle_results_addr "shuffle_results.addr"
               OpName %tid "tid"
               OpName %value "value"
               OpName %new_value "new_value"
               OpName %call "call"
               OpName %call1 "call1"
               OpName %call2 "call2"
               OpName %arrayidx "arrayidx"
               OpDecorate %33 Constant
         %33 = OpDecorationGroup
               OpDecorate %34 Alignment 4
         %34 = OpDecorationGroup
               OpDecorate %35 Alignment 8
         %35 = OpDecorationGroup
               OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
               OpDecorate %__spirv_BuiltInSubgroupLocalInvocationId BuiltIn SubgroupLocalInvocationId
               OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
               OpDecorate %__spirv_BuiltInSubgroupLocalInvocationId LinkageAttributes "__spirv_BuiltInSubgroupLocalInvocationId" Import
               OpGroupDecorate %33 %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInSubgroupLocalInvocationId
               OpGroupDecorate %34 %value %new_value
               OpGroupDecorate %35 %shuffle_results_addr %tid
      %ulong = OpTypeInt 64 0
       %uint = OpTypeInt 32 0
     %uint_1 = OpConstant %uint 1
    %v3ulong = OpTypeVector %ulong 3
%_ptr_UniformConstant_v3ulong = OpTypePointer UniformConstant %v3ulong
%_ptr_UniformConstant_uint = OpTypePointer UniformConstant %uint
       %void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
         %11 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint
%_ptr_Function__ptr_CrossWorkgroup_uint = OpTypePointer Function %_ptr_CrossWorkgroup_uint
%_ptr_Function_ulong = OpTypePointer Function %ulong
%_ptr_Function_uint = OpTypePointer Function %uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_UniformConstant_v3ulong UniformConstant
%__spirv_BuiltInSubgroupLocalInvocationId = OpVariable %_ptr_UniformConstant_uint UniformConstant
         %12 = OpFunction %void DontInline %11
%shuffle_results = OpFunctionParameter %_ptr_CrossWorkgroup_uint
         %14 = OpLabel
%shuffle_results_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function
        %tid = OpVariable %_ptr_Function_ulong Function
      %value = OpVariable %_ptr_Function_uint Function
  %new_value = OpVariable %_ptr_Function_uint Function
               OpStore %shuffle_results_addr %shuffle_results Aligned 8
         %22 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId
       %call = OpCompositeExtract %ulong %22 0
               OpStore %tid %call Aligned 8
      %call1 = OpLoad %uint %__spirv_BuiltInSubgroupLocalInvocationId
               OpStore %value %call1 Aligned 4
         %25 = OpLoad %uint %value Aligned 4
      %call2 = OpSubgroupShuffleINTEL %uint %25 %uint_1
               OpStore %new_value %call2 Aligned 4
         %28 = OpLoad %uint %new_value Aligned 4
         %29 = OpLoad %_ptr_CrossWorkgroup_uint %shuffle_results_addr Aligned 8
         %30 = OpLoad %ulong %tid Aligned 8
   %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %29 %30
               OpStore %arrayidx %28 Aligned 4
               OpReturn
               OpFunctionEnd

您可以看到，不方便进行读取或分析，我们将稍后讨论，这实际上是一个优势。

SPIR-V 为 OpenCL™ 应用开发人员提供的优势

白皮书中列举的许多优势都针对英特尔等驱动程序厂商，但是我们认为有必要考虑它为 OpenCL 应用开发人员提供的优势：

一旦生成 SPIR-V 内核后，其编译速度比文本内核 (OpenCL C) 快得多。无需在用户设备上执行基于文本的解析，因此，您可以更快速地启动应用。

应用未配备受保护的知识产权。如果配备了 OpenCL C 文本内核，更容易确定需要使用的算法。内核用 SPIR-V 编写时，需要反向工程。

SPIR-V 更接近硬件。以 OpenCL C 级别不可能实现的方式微调应用的性能。

它可以轻松扩展。研究人员或编译器发烧友可能想要使用开源英特尔® 显卡编译器和扩展的 SPIR-V 代码测试加速应用的新方法。

Khronos 积极开发了 SPIR-V 规范和生态系统。GitHub* 上提供了面向验证、汇编、反汇编、转换和优化库的工具，如 SPIRV-Tools、SPIRV-Cross、SPIRV-LLVM-Translator。

英特尔® 显卡驱动程序中的 SPIR-V*

我们来看一下如何在英特尔® 显卡驱动程序中编译 OpenCL 内核。在 OpenCL 2.1 中，可以通过多种方法将来自内核的源代码提供给驱动程序：

ClCreateProgramWithBinary
该库完全依赖于创建它的设备，应用不创建它。这意味着二进制文件不可移植到其他平台。这可能主要用于应用缓存内核。

您也可以使用 cl_khr_spir 扩展来提供 SPIR 格式的内核。SPIR 标准支持最高 2.0 的 OpenCL，在未来版本中，将替换为 SPIR-V。

ClCreateProgramWithSource
这是最常见的路径。内核源为文本格式，使用 OpenCL C 高级语言编写内核。英特尔® 显卡驱动程序中的编译流程如下所示：

图1.对使用 OpenCL C 编写的内核进行编译

在英特尔® 显卡编译器中，需要调用 Clang 将 OpenCL C 内核转换为 SPIR-V。然后，英特尔® 显卡编译器 OpenCL 前端将其转换为 LLVM-IR，编译器用它在内部生成硬件指令（机器代码）。

ClCreateProgramWithIL
该调用提供实施定义的中间语言 (IL)，它不能像在 ClCreateProgramWithBinary 中那样移植或使用 SPIR-V（Khronos 定义的可移植 IL）。
您可以看到，交付应用时准备 SPIR-V 可节省客户机上的编译时间，此处不需要 Clang 阶段：

图2.对使用 SPIR-V* 编写的内核进行编译

如何为现有的 OpenCL 内核生成 SPIR-V

如前所示，在驱动程序堆栈中，SPIR-V 由 Clang 生成。Clang 的独立版可从 Khronos Group GitHub* 中下载。

成功构建后，生成 SPIR-V 很简单：

clang -cc1 -emit-spirv -triple=spir64-unknown-unknown -cl-std=CL2.0 -include opencl.h kernel.cl -o kernel.spir

然后使用应用中带 clCreateProgramWithIL OpenCL API 调用的内核。

如何手动微调 SPIR-V* 内核

SPIR-V 是一种二进制格式，因此需要汇编程序和反汇编程序。

您可以从 KhronosGroup/SPIRV-Tools 中下载：

基本步骤为：

1. 使用 spirv-dis对 SPIR-V 内核进行反汇编
2. 编辑反汇编
3. 使用 spirv-as 再次汇编
4. 使用 spirv-val 进行验证

结论

自从推出 OpenCL 2.1 和 Vulkan 1.0 API 以来，英特尔® 显卡编译器便支持 SPIR-V。在 Vulkan 中，SPIR-V 是我们支持的唯一着色器表示。在 OpenCL 2.1 中，使用 OpenCL C 或 SPIR-V。

大多数 OpenCL 应用使用 OpenCL C，如前所示，需要转换为中间表示。最近，我们将该任务从 SPIR 迁移到 SPIR-V。这是一种更全面地测试与验证 SPIR-V 的方法，展示了我们支持 Khronos 标准的承诺。

资源

• 英特尔® 显卡编译器
• SPIRV-Tools
• SPIRV-Cross
• SPIRV-LLVM-Translator

脚注

1. Khronos SPIR-V 注册表
2. Khronos SPIR 注册表