diff --git a/YaeBlog/source/drafts/heterogeneous-programming-model.md b/YaeBlog/source/drafts/heterogeneous-programming-model.md index d21e86b..fcd17ec 100644 --- a/YaeBlog/source/drafts/heterogeneous-programming-model.md +++ b/YaeBlog/source/drafts/heterogeneous-programming-model.md @@ -290,8 +290,8 @@ std::vector> omp_matrix_multiply( | 运行方法 | 运行时间 | 比率 | | ------------ | -------- | ---- | -| SingleThread | 76823 ms | 1.00 | -| OpenMP | 8324 ms | 0.10 | +| SingleThread | 21685 ms | 1.00 | +| OpenMP | 2268 ms | 0.10 | ### CUDA @@ -426,10 +426,10 @@ std::vector> cudaCalculateMatrix(const std::vector selectDevice() +{ + cl_uint platformCount; + checkOpenCLError(clGetPlatformIDs(0, nullptr, &platformCount)); + std::cout << "Platform count: " << platformCount << std::endl; + + std::vector platforms(platformCount); + checkOpenCLError(clGetPlatformIDs(platformCount, platforms.data(), nullptr)); + + std::unique_ptr selectedDevice = nullptr; + + for (const auto& platform : platforms) + { + cl_uint deviceCount = 0; + checkOpenCLError(clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, 0, nullptr, &deviceCount)); + + std::vector devices(deviceCount); + checkOpenCLError(clGetDeviceIDs(platform, CL_DEVICE_TYPE_ALL, deviceCount, devices.data(), nullptr)); + + for (const auto& device : devices) + { + size_t deviceNameLength; + checkOpenCLError(clGetDeviceInfo(device, CL_DEVICE_NAME, 0, nullptr, &deviceNameLength)); + + std::vector deviceNameArray(deviceNameLength); + checkOpenCLError( + clGetDeviceInfo(device, CL_DEVICE_NAME, deviceNameLength, deviceNameArray.data(), nullptr)); + + std::string deviceName(deviceNameArray.data(), deviceNameArray.size() - 1); + + std::cout << "Found device: " << deviceName << std::endl; + + if (deviceName.find("4060") != std::string::npos) + { + std::cout << "Select device '" << deviceName << "' as runner." << std::endl; + selectedDevice = std::make_unique(); + selectedDevice->platform = platform; + selectedDevice->device = device; + } + else + { + clReleaseDevice(device); + } + } + } + + if (selectedDevice == nullptr) + { + std::cout << "Failed to find the target device." << std::endl; + std::exit(EXIT_FAILURE); + } + + return selectedDevice; +} + +std::vector clCalculateMatrix(const std::vector& a, + const std::vector& b) +{ + cl_int error; + + const std::unique_ptr computationContext = selectDevice(); + // A key-value list ends with 0 + // See also https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_API.html#context-properties-table + std::array properties = { + CL_CONTEXT_PLATFORM, + reinterpret_cast(computationContext->platform), + 0 + }; + + cl_context context = clCreateContext(properties.data(), 1, &computationContext->device, nullptr, nullptr, + &error); + checkOpenCLError(error); + cl_command_queue queue = clCreateCommandQueueWithProperties(context, computationContext->device, nullptr, + &error); + checkOpenCLError(error); + + std::vector result(MATRIX_SIZE * MATRIX_SIZE, 0); + constexpr size_t matrixSize = MATRIX_SIZE * MATRIX_SIZE * sizeof(int); + + cl_mem deviceA = clCreateBuffer(context, CL_MEM_READ_ONLY, matrixSize, nullptr, &error); + checkOpenCLError(error); + cl_mem deviceB = clCreateBuffer(context, CL_MEM_READ_ONLY, matrixSize, nullptr, &error); + checkOpenCLError(error); + cl_mem deviceC = clCreateBuffer(context, CL_MEM_READ_WRITE, matrixSize, nullptr, &error); + checkOpenCLError(error); + + checkOpenCLError( + clEnqueueWriteBuffer(queue, deviceA, CL_TRUE, 0, matrixSize, a.data(), 0, nullptr, + nullptr)); + checkOpenCLError( + clEnqueueWriteBuffer(queue, deviceB, CL_TRUE, 0, matrixSize, b.data(), 0, nullptr, + nullptr)); + // Copy result to erase the previous result + checkOpenCLError( + clEnqueueWriteBuffer(queue, deviceC, CL_TRUE, 0, matrixSize, result.data(), 0, + nullptr, nullptr + )); + + auto source = R"( +#define MATRIX_SIZE 2048 + +__kernel void calculate(const __global int* a, const __global int* b, __global int* c) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + + int result = 0; + for (int i = 0; i < MATRIX_SIZE; i++) + { + result += a[x * MATRIX_SIZE + i] * b[i * MATRIX_SIZE + y]; + } + + c[x * MATRIX_SIZE + y] = result; +})"; + + cl_program program = clCreateProgramWithSource(context, 1, &source, nullptr, &error); + checkOpenCLError(error); + checkOpenCLError(clBuildProgram(program, 0, nullptr, "", nullptr, nullptr)); + + size_t messageSize; + checkOpenCLError( + clGetProgramBuildInfo(program, computationContext->device, CL_PROGRAM_BUILD_LOG, 0, nullptr, &messageSize)); + std::vector messageArray(messageSize); + checkOpenCLError( + clGetProgramBuildInfo(program, computationContext->device, CL_PROGRAM_BUILD_LOG, messageSize, messageArray.data( + ), nullptr)); + std::string message(messageArray.data(), messageSize - 1); + std::cout << "Build log: " << message << std::endl; + + cl_kernel kernel = clCreateKernel(program, "calculate", &error); + checkOpenCLError(error); + + checkOpenCLError(clSetKernelArg(kernel, 0, sizeof(cl_mem), &deviceA)); + checkOpenCLError(clSetKernelArg(kernel, 1, sizeof(cl_mem), &deviceB)); + checkOpenCLError(clSetKernelArg(kernel, 2, sizeof(cl_mem), &deviceC)); + + cl_event event; + constexpr std::size_t globalSize[2] = {MATRIX_SIZE, MATRIX_SIZE}; + checkOpenCLError(clEnqueueNDRangeKernel(queue, kernel, 2, nullptr, + globalSize, nullptr, 0, nullptr, &event)); + + checkOpenCLError(clWaitForEvents(1, &event)); + + checkOpenCLError( + clEnqueueReadBuffer(queue, deviceC, CL_TRUE, 0, matrixSize, result.data(), 0, + nullptr, nullptr)); + + clReleaseMemObject(deviceA); + clReleaseMemObject(deviceB); + clReleaseMemObject(deviceC); + + clReleaseKernel(kernel); + clReleaseProgram(program); + clReleaseCommandQueue(queue); + clReleaseContext(context); + clReleaseDevice(computationContext->device); + return result; +} +``` + +从上面的代码中可以看出两点： + +- OpenCL的编程比CUDA的更为繁琐，因为OpenCL支持的设备种类更多，在主机代码上还需要多出一块选择运行设备的代码； +- OpenCL在主机代码和核函数的解耦更为彻底，核函数直接以字符串的形式存在于主机代码中，而各个厂商提供的驱动才是真正的编译器。 + +测试的运行结果如下： + +| 类型 | 运行时间 | 比率 | +| ----------------------------- | -------- | ---- | +| NVIDIA 4060 Ti OpenCL | 173ms | 0.01 | +| Intel UHD Graphics 770 OpenCL | 1020ms | 0.04 | +| CPU | 21255ms | 1.00 | ### SYCL +SYCL是一个使用标准C++编写在各种异构计算设备上运行核函数的抽象层，并提供了一套新的API来查找各种设备并管理这些设备上的内存资源和代码执行。这个标准是开发、无版税、跨平台的抽象标准。同时也是因为这是一个**标准**，因此需要寻找支持这个标准的编译器才能使用这个标准。按照官网上的说明，我们选择了两个看上去还在活跃开发的项目，Intel的[oneAPI](https://www.intel.com/content/www/us/en/developer/tools/oneapi/overview.html)和开源的[AdaptiveCpp](https://github.com/AdaptiveCpp/AdaptiveCpp)进行调研，考虑到在后文中还将继续介绍oneAPI相关的工作，因此这里将重点放在AdaptiveCpp上。 +AdaptiveCpp由四个部分组成，分别在不同的C++命名空间中提供。 -### OpenAcc +- SYCL Interface：实现了SYCL标准中规定的各种类和函数，是实际上同用户交互的接口。这些接口实际上可以仍然可以分成主机API和核函数库两个部分。主机API是普通的C++代码，负责任务调度、任务管理和平台射别管理等。核函数库包括了这种在编写核函数时可以使用的类和函数，这些接口暴露一些后端特定的功能，其中的一些甚至需要使用后端特定的方言来编写，例如CUDA。 +- AdaptiveCpp Runtime：运行时实际上实现了设备调度、任务图管理和执行、数据管理、后端管理、任务调度和同步等等功能，运行时负责同各种支持后端的运行时交互来实现上述的功能。 + ![image-20241029123308139](./heterogeneous-programming-model/image-20241029123308139.png) -### Triton +- Compiler：考虑到在用户编写的代码中可能使用一些特定后端的方言，因此普通的C++编译器无法正常编译所有的用户代码。因此用户代码的编译是通过一个名为`acpp`的Python脚本驱动的，这个脚本将各个后端的不同编译器暴露为一个统一的编程接口。 + +- Glue：将上述的各个部分连接在一起的胶水代码。一种典型的胶水代码是内核函数的启动代码`kernel launcher`，由于启动器中往往涉及到一些后端特定的方言，例如CUDA中的`<<<>>>`或者OpenMP中的各种`pragma`，因此这些代码通常需要使用特定的编译器进行编译，所以这些胶水代码直接以头文件的方式提供，以方便在编译时被特定的编译器处理。这些胶水代码将会把核函数包裹为一个合法的C++函数对象，这样运行时就可以获得这个函数对象并控制代码在设备上的运行。 + +AdaptiveCpp同时支持多种不同的编译流程。 + +1. 一种通用的一遍编译流程，将核函数编译到一种统一的中间表示形式，这种中间表示形式将在运行时被编译到特定的后端架构上。这种编译流程提供了高度的可移植性和较快的编译速度。这种编译设施支持的后端有：通过`PTX`在NVIDIA的GPU上运行，通过`amdgcn`在AMD的GPU上运行，通过`SPIR-V`在Intel的GPU上运行，通过`SPIR-V`在任何支持OpenCL驱动的设备上运行，也可以通过LLVM直接在CPU上运行。 +2. 一种为互操作性优化的多遍编译流程，在这个流程中AdaptiveCpp将聚合现有的各种LLVM/Clang的编译工具链，使得用户可以在单个代码文件中混合编写SYCL和各种特定的编程模型，例如CUDA和HIP。使用这个编译流程的好处有亮点：（1）在这种编译流程中可以直接在SYCL代码使用各个特定编译模型中提供最新设备内部优化（Intrinsics），不用等待SYCL标准的支持；（2）在这种编译流程中可以使用各个厂商提供的优化模板库，例如`rocPRIM`和`CUB`。这种编译流程是提供聚合`CUDA`的clang前端和`ROCm`的clang前端来实现的。 +3. 一种只将AdaptiveCpp作为函数使用的编程流程。在这种情况AdaptiveCpp作为一个三方库被引入其他的编译器编译流程中。 + +第一种通用的编译流程显然是泛用性最广的一种编译流程，同时也是AdaptiveCpp推荐的编译流程。 + +![image-20241029163654675](./heterogeneous-programming-model/image-20241029163654675.png) + +下面是一段使用SYCL进行矩阵乘法加速的代码： + +```cpp +struct CustomDeviceSelector +{ + explicit CustomDeviceSelector(std::string vendorName) : _vendorName(std::move(vendorName)) + { + } + + int operator()(const sycl::device& d) const + { + int deviceRating = 0; + + if (d.is_gpu() && d.get_info().find(_vendorName) != std::string::npos) + { + deviceRating = 3; + } + else if (d.is_cpu()) + { + deviceRating = 1; + } + + return deviceRating; + } + +private: + std::string _vendorName; +}; + +static std::vector syclCalculateMatrix(const std::vector& a, const std::vector& b, + const std::string& hint) +{ + const CustomDeviceSelector selector(hint); + sycl::queue queue(selector); + + const std::string deviceName = queue.get_device().get_info(); + std::cout << "Select device: " << deviceName << std::endl; + + std::vector result(MATRIX_SIZE * MATRIX_SIZE, 0); + + sycl::buffer aBuffer(a); + sycl::buffer bBuffer(b); + sycl::buffer resultBuffer(result); + + queue.submit([&](sycl::handler& h) + { + const sycl::accessor aBufferAccessor(aBuffer, h, sycl::read_only); + const sycl::accessor bBufferAccessor(bBuffer, h, sycl::read_only); + const sycl::accessor resultBufferAccessor(resultBuffer, h, sycl::write_only); + + h.parallel_for(sycl::nd_range<2>({MATRIX_SIZE, MATRIX_SIZE}, {16, 16}), [=](const sycl::nd_item<2>& item) + { + const size_t x = item.get_global_id(0); + const size_t y = item.get_global_id(1); + + int temp = 0; + for (size_t k = 0; k < MATRIX_SIZE; ++k) + { + temp += aBufferAccessor[x * MATRIX_SIZE + k] * bBufferAccessor[k * MATRIX_SIZE + y]; + } + resultBufferAccessor[x * MATRIX_SIZE + y] = temp; + }); + }); + + sycl::host_accessor resultHostAccessor(resultBuffer, sycl::read_only); + + for (size_t i = 0; i < MATRIX_SIZE; ++i) + { + for (size_t j = 0; j < MATRIX_SIZE; ++j) + { + result[i * MATRIX_SIZE + j] = resultHostAccessor[i * MATRIX_SIZE + j]; + } + } + + return result; +} +``` + +测试之后的运行结果如下所示： + +| 类型 | 运行时间 | 比率 | +| --------------------------- | -------- | ----- | +| Intel UHD Graphics 770 SYCL | 488ms | 0.023 | +| NVIDIA 4060 Ti SYCL | 180ms | 0.008 | +| OpenMP SYCL | 1591ms | 0.076 | +| CPU | 20930ms | 1.000 | + +### OpenACC + +OpenACC是一个通过编译器制导来在代码中表达并行性并利用并行编译器为多个并行加速器生成代码的编程模型。为了保证OpenACC可以适配于各种计算架构的加速设备，OpenACC设计了一个各种并行层次和有着不同速度和寻址方式内存的编程模型。同时OpenACC主要的功能即是支持同时将计算和数据卸载到一个加速设备上，考虑到加速设备可能有着同宿主设备完全不同的内存架构，OpenACC编译器和运行时将会自动分析代码并负责加速器上内存的管理和加速器和主机之间的数据传输。 + +作为一个高等级、平台独立的加速器编程框架，使用OpenACC进行开发能够使开发人员将一个源代码编译到一系列设备上运行并实现一个相对较好的性能，但是这个简易性和移植性也在一定程度上造成使用OpenACC编程无法完全利用加速设备上的算力。 + +OpenACC是作为一个标准的形式提供的，实现了该标准的编译器有： + +| 编译器名称 | 情况 | +| ------------------------------------------------------------ | ------------------------------------------------------------ | +| NVIDIA HPC SDK | 支持在NVIDIA GPU和多核CPU上的OpenACC并行编程 | +| Sourcery CodeBench Lite | OpenACC官网上说支持针对AMD GPU的编译，但是官网页面似乎改版了，没有找到相关的内容 | +| GCC 12 | 支持到OpenACC 2.6 | +| [Omni Compiler Project](https://github.com/omni-compiler/omni-compiler) | 源到源编译器，将带有制导的源代码翻译到带有运行时调用的平台代码，近两年没有活跃开发 | +| [OpenUH](https://github.com/uhhpctools/openuh) | 项目开发者在7年前的最后一次提交了中删除了README中有关OpenACC的内容 | @@ -486,11 +783,17 @@ OpenCL是目前最为典型、发展最好的异构并行编程模型，毕竟 +### Triton + ## 参考文献 -1. 刘颖,吕方,王蕾,陈莉,崔慧敏,冯晓兵.异构并行编程模型研究与进展.软件学报,2014,25(7):1459-1475. http://www.jos.org.cn/1000-9825/4608.htm -2. +1. 刘颖,吕方,王蕾,陈莉,崔慧敏,冯晓兵.异构并行编程模型研究与进展.软件学报,2014,25(7):1459-1475. [http://www.jos.org.cn/1000-9825/4608.htm](http://www.jos.org.cn/1000-9825/4608.htm) +2. AdaptiveCpp官方文档. [https://adaptivecpp.github.io/AdaptiveCpp/](https://adaptivecpp.github.io/AdaptiveCpp/) +3. Exploring the performance of SGEMM in OpenCL on NVIDIA GPUs. [https://github.com/CNugteren/myGEMM](https://github.com/CNugteren/myGEMM) +4. OpenACC Programming and Best Practices Guide. [https://openacc-best-practices-guide.readthedocs.io/en/latest/01-Introduction.html](https://openacc-best-practices-guide.readthedocs.io/en/latest/01-Introduction.html) +5. oneAPI What is it?. [https://www.intel.com/content/www/us/en/developer/articles/technical/oneapi-what-is-it.html](https://www.intel.com/content/www/us/en/developer/articles/technical/oneapi-what-is-it.html) +6. diff --git a/YaeBlog/source/drafts/heterogeneous-programming-model/image-20241029123308139.png b/YaeBlog/source/drafts/heterogeneous-programming-model/image-20241029123308139.png new file mode 100644 index 0000000..774102e --- /dev/null +++ b/YaeBlog/source/drafts/heterogeneous-programming-model/image-20241029123308139.png @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:9675685943c6124ffdc70023f19c8d043c484d72f62866e6b4e87eca270e0366 +size 76309 diff --git a/YaeBlog/source/drafts/heterogeneous-programming-model/image-20241029163654675.png b/YaeBlog/source/drafts/heterogeneous-programming-model/image-20241029163654675.png new file mode 100644 index 0000000..2d64a1a --- /dev/null +++ b/YaeBlog/source/drafts/heterogeneous-programming-model/image-20241029163654675.png @@ -0,0 +1,3 @@ +version https://git-lfs.github.com/spec/v1 +oid sha256:4e38a5bdee31a5ddf11329d6b72ca8b1b2b119f5988166477e8e0ff63a1fa40d +size 195933