DPC++简介

Ⅲ 初探DPC++
在开始讨论现代C++语言在DPC++中的应用之前，让我们先看一遍完整的代码，顺便测试我们的实验环境：

#include <CL/sycl.hpp>
constexpr int N = 16;
using namespace sycl;

class IntelGPUSelector : public device_selector {
public:
int operator()(const device& Device) const override {
const std::string DeviceName = Device.get_infoinfo::device::name();
const std::string DeviceVendor = Device.get_infoinfo::device::vendor();

return Device.is_gpu() && (DeviceName.find("Intel") != std::string::npos) ? 100 : 0;
1

}
};

int main() {
IntelGPUSelector d;
queue q(d);
int* data = malloc_shared(N, q);
q.parallel_for(N, [=](auto i) {
data[i] = i;
}).wait();
for (int i = 0; i < N; i++) std::cout << data[i] << " ";
free(data, q);
}
编译运行上面的代码，如果没有问题应该输出：

0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15
简单解释一下这段代码，sycl是DPC++的实体的命名空间，用using namespace sycl;打开命名空间可以简化后续代码。IntelGPUSelector是一个继承了device_selector的设备选择器，其中device_selector是纯虚类，它有个纯虚函数int operator()(const device& Device) const需要派生类来实现，该函数会遍历计算机上的计算设备，并且返回使用设备的优先级，返回数字越高优先级越高，这里选择Intel的GPU作为首选的计算设备，注意这个函数使用了override来说明其目的是覆盖虚函数。queue的目的是指定工作的目标位置，这里设置的是Intel的GPU。函数模板malloc_shared分配了可在设备上使用的工作内存。成员函数parallel_for执行并行计算。值得注意的是free调用的是sycl::free而不是C运行时库的free。在这段代码中，比较明显使用了现在C++语法的地方是函数parallel_for的实参，

[=](auto i) { data[i] = i; }
这是一个lambda表达式。
Ⅳ DPC++和lambda表达式
如果要选出一个对DPC++最重要的现代C++语言特性，我觉得lambda表达式应该可以被选上。因为在DPC++的代码中，内核代码一般都是以lambda表达式的形式出现。比如上面的例子就是将lambda表达式作为对象传入到Intel的GPU设备上然后进行计算的。在这个lambda表达式中，[=]是捕获列表，它可以捕获当前定义作用域内的变量的值，这也是它可以在函数体内使用data[i]的原因。捕获列表[=]之后的是形参列表(auto i)，注意这里的形参类型使用的是auto占位符，也就是说，我们将形参类型的确认工作交给了编译器。我们一般称这种lambda表达式为泛型lambda表达式。当然，如果在编译时选择C++20标准，我们还可以将其改为模板语法的泛型lambda表达式：

[=](T i) { data[i] = i; }
lambda表达式的捕获列表功能非常强大，除了捕获值以外，还可以捕获引用，例如：

[&](auto i) { data[i] = i; }
以上代码会捕获当前定义作用域内的变量的引用，不过值得注意的是，由于这里的代码会交给加速核心运行，捕获引用并不是一个正确的做法，会导致编译出错。另外一般来说，我们并不推荐直接捕获所有可捕获的对象，而是有选择的捕获，例如：

[data](auto i) { data[i] = i; }
当然，除了使用lambda表达式，我们也可以选择其他形式的代码来运行设备，比如使用仿函数：

struct AssginTest {
void operator()(auto i) const { data_[i] = i; }
int* data_;
};

AssginTest functor{data};
q.parallel_for(N, functor).wait();
但是很明显，这种方法没有使用lambda表达式来的简单直接。

Ⅴ DPC++和泛型能力
之所以能够让parallel_for这么灵活的接受各种形式的实参，是因为parallel_for本身是一个成员函数模板：

template
event parallel_for(range<1> NumWorkItems,
_KERNELFUNCPARAM(KernelFunc) _CODELOCPARAM(&CodeLoc)) {
_CODELOCARG(&CodeLoc);
return parallel_for_impl(NumWorkItems, KernelFunc, CodeLoc);
}
其中KernelFunc就是传入的lambda表达式或者仿函数，KernelType是KernelFunc的类型。
如果从这里的代码一路运行跟踪下去，会发现它们都是用模板传递实参类型，直到submit_impl：

sycld.dll!cl::sycl::queue::submit_impl
dpcpp.exe!cl::sycl::queue::submit
dpcpp.exe!cl::sycl::queue::parallel_for_impl
dpcpp.exe!cl::sycl::queue::parallel_for
这是因为sycld.dll是一个二进制模块，它无法以模板的形式提供代码，所有的类型必须确定下来，为了解决这个问题，cl::sycl::queue::submit_impl使用了std::function：

event submit_impl(function_class<void(handler &)> CGH,
const detail::code_location &CodeLoc);
函数模板cl::sycl::queue::parallel_for_impl将KernelFunc封装到另外一个lambda表达式对象中，并且通过function_class<void(handler &)>来传递整个lambda表达式：

template <typename KernelName = detail::auto_name, typename KernelType,
int Dims>
event parallel_for_impl(
range NumWorkItems, KernelType KernelFunc,
const detail::code_location &CodeLoc = detail::code_location::current()) {
return submit(
[&](handler &CGH) {
CGH.template parallel_for<KernelName, KernelType>(NumWorkItems,
KernelFunc);
},
CodeLoc);
}
其中function_class就是std::function。注意这里CGH.template parallel_for需要说明符template否则尖括号会解析出错。DPC++通过这样一系列的操作，最大限度的保留了用户编程的灵活性。

Ⅵ DPC++和模板推导
DPC++代码中大量的运用了C++17标准才引入的模板推导特性，关于这些特性我们还是从一个DPC++的小例子开始：

int main() {
IntelGPUSelector d;
queue q(d);
std::vector v1(N);
std::array<int, N> v2;
{
buffer buf1(v1);
buffer buf2(v2);

q.submit([&](handler& h) {
  accessor a1(buf1, h, write_only);
  accessor a2(buf2, h, write_only);
  h.parallel_for(N, [=](auto i) {
    a1[i] = i;
    a2[i] = i;
  });
});
1
2
3
4
5
6
7
8

}
for (int i = 0; i < N; i++) std::cout << v1[i] << v2[i] << " ";
}
这段代码没有使用malloc_shared分配内存，取而代之的是使用buffer和accessor，其中buffer用于封装数据，accessor用于访问数据。这里以buffer为例解析DPC++对模板推导的使用。
首先观察buffer的两个实例，它们的构造函数的实参分别是std::vector和std::array<int, N>类型。之所以能够这样调用构造函数，并不是因为buffer为这两个类型重载了它的构造函数，而是因为其构造函数使用了模板。这里涉及到一个C++17标准新特性——类模板的模板实参推导。在以往，类模板的实例化必须是显式传入模板实参，否则会造成编译出错。在新的标准中，类模板的模板实参已经可以根据构造函数来推导了。来看一下buffer的构造函数：

template <typename T, int dimensions = 1,
typename AllocatorT = cl::sycl::buffer_allocator,
typename = typename detail::enable_if_t<(dimensions > 0) &&
(dimensions <= 3)>>
class buffer {
public:
…
template <class Container, int N = dimensions,
typename = EnableIfOneDimension,
typename = EnableIfContiguous>
buffer(Container &container, AllocatorT allocator,
const property_list &propList = {})
: Range(range<1>(container.size())) {
impl = std::make_shareddetail::buffer_impl(
container.data(), get_count() * sizeof(T),
detail::getNextPowerOfTwo(sizeof(T)), propList,
make_unique_ptr<detail::SYCLMemObjAllocatorHolder>(
allocator));
}

template <class Container, int N = dimensions,
typename = EnableIfOneDimension,
typename = EnableIfContiguous>
buffer(Container &container, const property_list &propList = {})
: buffer(container, {}, propList) {}
…
};
代码buffer buf1(v1);会执行

buffer(Container &container, const property_list &propList = {})
这条构造函数，值得注意的是该构造函数并没有实际的实现代码，而是通过委托构造函数的方法调用了

buffer(Container &container, AllocatorT allocator, const property_list &propList = {})
委托构造函数是C++11引入的特性，它可以让某个构造函数将构造的执行权交给另外的构造函数。回到模板推导，这里通过构造函数会推导出Container是std::vector，dimensions的推导结果是1，而后面两个模板参数是用来检查前两个模板参数是否正确的，这里大量的使用了模板元编程的技巧：

template
using EnableIfOneDimension = typename detail::enable_if_t<1 == dims>;

template
using EnableIfContiguous =
detail::void_t<detail::enable_if_t<std::is_convertible<
detail::remove_pointer_t<decltype(
std::declval().data())> ()[],
const T ()[]>::value>,
decltype(std::declval().size())>;
首先它们都是使用using定义的别名模板，它们的目的分别是检查dims是否为1和Container是否为连续的。第一个别名模板很简单，直接检查dims是否为1，detail::enable_if_t就是std::enable_if_t。第二个检查连续性的方法稍微麻烦一些，简单来说就是检查容器对象的成员函数data()返回值的类型的数组指针是否能和const T ()[]转换，这里主要检查两点，第一容器具有data()成员函数，第二返回类型的指针和T const T ()[]转换。事实上，在标准容器中，只有连续容器有data()成员函数，其他的都会因为没有data()而报错，例如：

no member named ‘data’ in ‘std::list’
仔细阅读上面代码的朋友应该会发现另外一个问题，那就是没有任何地方可以帮助编译器推导出buffer的类模板形参T。这就不得不说DPC++将C++17关于模板推导的新特性用的淋漓尽致了。实际上在代码中，有这样一句用户自定义推导指引的代码：

template
buffer(Container &, const property_list & = {})
->buffer<typename Container::value_type, 1>;
用户自定义推导指引是指程序员可以指导编译器如何通过函数实参推导模板形参的类型。最后在这个例子中，需要注意一下，buffer在析构的时候才会将缓存的数据写到v1和v2，所以这里用了单独的作用域。

~buffer_impl() {
try {
BaseT::updateHostMemory();
} catch (…) {
}
}
亚马逊测评 www.yisuping.cn