OpenCL SVM特性|极客笔记

SVM缓冲有多个特性，用来匹配不同的硬件支持和软件控制。并不要求硬件对这些特性都支持，可以只支持所有特性中的某些特性。每个特性可以单独考虑，不同的特性支持的功能也不一样。下面就来详细讲解SVM的这些特性。

共享虚拟地址空间

OpenCL 2.0中一个显著的特性就是共享虚拟内存，这意味着在主机上分配的指针，设备端内核也可以使用，反之亦然。但是该功能仅限于指向SVM缓冲的指针，而不是OpenCL中常规的缓冲区对象。对于设备端，只有指向全局内存的指针可以用来共享数据，如图5-12所示。
需要注意的是，共享虚拟内存和共享物理内存的区别。共享物理内存是主机和设备共享相同的物理内存，但是它们的虚拟地址可能并不相同。共享物理内存的好处是主机与设备之间可以有效地传输数据。共享物理内存在OpenCL标准中并未做说明，这依赖于具体的平台和设备（如AMD APU平台，ARM CPU与Mali GPU）。而共享虚拟内存只是软件功能的扩展，在实际硬件环境中主机与设备可能有各自独立的存储空间，也可能是共享物理内存的。
SVM缓冲共享虚拟地址空间特性，适用于粗粒度和细粒度SVM。

无须映射访问

对于主机与设备间的交互，一个很重要的机制就是映射/解映射OpenCL缓冲区域。当底层硬件无法做到主机与设备细粒度访问OpenCL缓冲时，就需要映射/解映射OpenCL缓冲区域。映射/解映射OpenCL缓冲需要使用OpenCL API函数显式操作。当主机与设备间真正需要细粒度数据交换时，映射/解映射OpenCL缓冲显得太冗长与烦琐。
随着硬件技术发展，OpenCL平台可以摆脱显式的映射/解映射命令。这种情况使得主机与设备可以以任何粒度访问SVM缓冲，把保持内存一致性的问题交给底层硬件去处理。SVM无须映射访问的特性，只适用于细粒度SVM。下表是SVM不同粒度访问方式比较。

SVM不同粒度访问方式比较

对比如上粗粒度和细粒度SVM的处理过程代码，在每次主机与设备间交互时，无须映射/解映射访问，带来了编程上的便利性。
为了创建无须映射访问的SVM缓冲，clSVMAlloc函数应该传入CL_MEM_SVM_FINE_GRAIN_BUFFER标志，如下代码：

void ＊p = clSVMAlloc(
                        context,             //一个有效的OpenCL上下文
                        CL_MEM_READ_WRITE | CL_MEM_SVM_FINE_GRAIN_BUFFER,
                        size,//分配的SVM缓冲大小,单位为字节
                        0                    //字节对齐(0采用默认值)
                        );

细粒度一致的访问

细粒度SVM缓冲提供主机和设备同时修改相同内存区域的访问。也就是说主机端使用clEnqueueNDRangeKernel命令入队一个内核，同时无须等待内核完成执行，主机可以修改相同的SVM缓冲。下表是细粒度一致访问示例。

细粒度一致访问示例

在clFinish（）以后，主机和设备对于存储器P中的对象能看到完全相同的内容，存储器对象的值为{0.f,1.f,2.f,3.f,4.f,5.f}。
主机和设备读写SVM缓冲的不同字节位置时，必须在OpenCL同步点保证内存一致性。在例子中主机和设备读写SVM缓冲不同字节位置，如果读写SVM缓冲相同字节位置，需要额外的同步点，如原子操作和内存栅栏。
SVM细粒度一致性访问的特性，只适用于细粒度SVM。

细粒度同步

SVM细粒度同步的特性，可以实现主机（或设备）对SVM缓冲的修改对设备（或主机）也有效，而不需要入队任何数据传输命令（如缓冲读或映射/解映射）。主机和设备在SVM缓冲相同位置执行并发原子操作，协作实现该功能。主机和设备也可以使用内存栅栏提供必要的内存一致性。
应用于SVM缓冲的原子操作具有以下属性，它们能够为主机与设备同步提供一种强有力的机制：

访问原子性：对标量类型的某个特定变量的事务（transactional）访问。例如，对于int类型，使用原子操作主机和设备都可以安全地更新特定的整数变量。
内存一致性：确保主机（或设备）对SVM缓冲的读写对设备（或主机）是可见的，并且前后顺序正确。例如，在SVM缓冲中有一个循环队列，往循环队列中插入一新的队列项，需要更新主机上队列的next_item指针变量，设备能够以正确的顺序看到指针变量的更新结果。为了提供这个支持，OpenCL 2.0使用原子操作时有多个可以明确指定的排序规则。

为了使用原子操作和栅栏，需要在内核中指定memory_scope_all_svm_devices内存范围，同时在创建SVM缓冲时需要指定CL_DEVICE_SVM_ATOMICS标志，如下：

void ＊p = clSVMAlloc(
                context,           //一个有效的OpenCL上下文
                CL_MEM_READ_WRITE
                | CL_MEM_SVM_FINE_GRAIN_BUFFER
                | CL_DEVICE_SVM_ATOMICS,
                size,//分配的SVM缓冲大小,单位为字节
                0                   //字节对齐(0采用默认值)
            );

SVM细粒度同步的特性，只适用于支持原子操作的粗粒度SVM。
如下代码展示了同时初始化一个浮点型数组。数组元素同时由主机和设备初始化。元素索引在SVM缓冲中作为原子递增计数器，在主机与设备间共享：
主机代码：

//作为共享原子递增计数器
auto index = (std::atomic〈cl_int> ＊)clSVMAlloc(...,
              CL_MEM_READ_WRITE |
              CL_MEM_SVM_FINE_GRAIN_BUFFER |
              CL_MEM_SVM_ATOMICS,
              sizeof(cl_int), 0);
//分配SVM缓冲,同时会被主机和设备写入
auto p = (float ＊)clSVMAlloc (...,
                                  CL_MEM_READ_WRITE |
                                  CL_MEM_SVM_FINE_GRAIN_BUFFER,
                                  size ＊sizeof(float), 0);
clSetKernelArgSVMPointer(kernel, 0, index);
clSetKernelArgSVMPointer(kernel, 1, p);
clEnqueueNDRangeKernel(..., kernel, ...
                        &size, ...);    //全局大小与p中元素个数相同
clFlush(...);
int localIndex;
while((localIndex =
          std::atomic_fetch_add_explicit(
                  index, 1,
                  std::memory_order_relaxed) ) 〈 size)
{
    p[localIndex] = localIndex;
}
clFinish(...);
//主机与设备同时以任意顺序初始化指针p。每个元素只初始化一次。

内核代码：

kernel void mykernel (global atomic_int ＊index, global float ＊p)
{
    int localIndex = atomic_fetch_add_explicit(
                        index, 1,
                        memory_order_relaxed,
                        memory_scope_all_svm_devices
                    );
    if(localIndex 〈 get_global_size(0))
      p[localIndex] = localIndex;
}

共享整个主机地址空间

OpenCL 2.0平台支持系统SVM，允许设备上的内核使用主机地址空间的任何数据。不需要调用clSVMAlloc来分配SVM内存。对于任何主机可用的内存（如malloc函数或new操作符分配的地址空间），设备上的内核都可以使用。OpenCL 2.0系统SVM地址空间示意图如下图所示。

OpenCL 2.0系统SVM地址空间示意图

系统SVM共享整个主机地址空间的特性对于没有内存分配权限的应用非常重要，如使用的库在内部分配内存。另一个例子就是把现有C/C++应用移植到OpenCL上，如果应用十分庞大和复杂，许多地方会分配内存。把这样的程序用OpenCL 2.0缓冲SVM移植工作量比较大，因为每个内存分配的代码都需要使用clSVMAlloc函数重写。而若使用系统SVM，则不需要重写代码。
尽管可以使用主机地址空间中的任何数据，但是还是需要根据OpenCL标准适当对齐数据。甚至于为了数据访问更有效，需要更强的对齐规则，这取决于使用的OpenCL平台。
SVM共享整个主机地址空间的特性，只适用于细粒度系统SVM。
下表展示了缓冲SVM和系统SVM用法区别。

缓冲SVN和系统SVM的用法区别

当OpenCL平台不支持系统SVM时，采用左边的方式去处理；只有当OpenCL平台支持系统SVM时，右边的方式才是正确。对于系统SVM，不需要使用clSVMAlloc函数来分配SVM缓冲在主机与设备间共享数据。

隐式地使用SVM缓冲

在OpenCL中，内核使用的任何缓冲都需要使用clSetKernelArg作为参数传入内核。对于不支持细粒度系统SVM的OpenCL平台，OpenCL 2.0 SVM缓冲也有类似的操作。对于SVM缓冲，调用如下函数把SVM缓冲作为参数传递给内核。
1）clSetKernelArgSVMPointer：把指向SVM缓冲的一个指针作为一个内核参数传递。
2）clSetKernelExecInfo：把指向内核使用的SVM缓冲的指针传递给内核，但不是作为内核参数传递。例如，一个指向SVM缓冲A的指针存储的是另一个SVM缓冲B，可以把缓冲B使用该函数传递给内核。
对于缓冲SVM，只能采用上述两种方式传递给内核，其他方式都被禁止的。在某些情况下，需要分配很多的缓冲SVM，但是OpenCL平台限制了每个内核使用的SVM缓冲个数。对于这种情况，应该如何处理呢？
对于细粒度系统SVM来说，不需要把SVM缓冲作为参数传递给内核，因为内核也可以访问SVM缓冲中的数据。因此，对于上述限制SVM缓冲个数的问题，可以使用细粒度系统SVM来解决。
SVM隐式地使用SVM缓冲的特性，只适用于细粒度系统SVM。
下表的例子展示了细粒度缓冲SVM和细粒度系统SVM的用法区别。

细粒度缓冲SVM和细粒度系统SVM用法区别

右边代码中，node2没有做任何改动，内核可以直接访问，极大地方便了程序开发。
当然如果内核不使用任何细粒度系统SVM，而只在内核中使用细粒度缓冲SVM，应用可以选择通知运行时，具体操作代码如下：

cl_bool flag = CL_FALSE;
clSetKernelExecInfo(mykernel,
                        CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM,
                        sizeof(flag),
                        &flag)