Intel OneAPI并行编程实验：以并行卷积为例

用户014295878856
249 阅读6分钟

OneAPI介绍

一种通用的可以适配Intel的多种异构硬件的并行编程库，个人理解可以用于HPC场景以及深度学习的卷积优化、向量传播等场景，并降低适配不同硬件的编译适配成本。

OneAPI提供了云平台DevCloud可以免费注册后使用，即使笔记本不是Intel的也可以使用Intel OneAPI配套硬件的体验，还无需自行安装toolkit。

jupyter.oneapi.devcloud.intel.com/ 这是我在本次实验中使用的云平台链接，本次实验主要基于Jupyter Lab进行，使用简单，用Jupyter Notebook构建了可视化界面，即使不会使用命令行进行编译也可以直接复制官方教程中的脚本进行编译，免去了CPP的编译痛苦。

问题描述

使用基于oneAPI的C++/SYCL实现一个用于计算图像的卷积操作。输⼊为一个图像矩阵和一个卷积核矩阵，输出为卷积后的图像。

实验过程

分析问题，图像的卷积是一个二维操作。

出于简便性，先考虑一维卷积的情况（例如在声音信号处理的场合），由于卷积结果的每一个位置的值都是一次计算即可求出的，所以卷积计算天生就是可以并行的，最理想的情况当然是对于长度为NN的信号块，使用NN个线程进行并行计算，然后将每次计算的结果写入对应的位置。

对于图像卷积而言，卷积核是二维的，按照上述思路，每一个目标像素点都由一个线程进行计算后写入对应位置。

可能的问题与优化：如果图像的大小是N2N^2，那么就需要同样多数量的线程进行计算，但，真的有这么多线程吗？

优化思路：将卷积也分成块状计算，例如对于一个K2K^2的卷积核，原始图像也分成K2K^2大小的块进行计算，这样（1）可以减少需要的线程数，更加现实；（2）可以利用cache的局部性原理，对原图像中存储在相近位置（一行）的数据多次重复访问，减少cache miss次数，这在原图像较大时可能很有意义。

代码

44行定义了提交的并行任务队列q，45~47行将buffer转化为只读或只写的accessor用于读取和写入，这是在语言自带的数据结构层面上的又一层抽象，保证了异构的计算设备的代码通用性。

49行使用二维的nd_item用于处理二维卷积图像，如果是一维的卷积，使用id<1>就可以了。

86行开始的验证部分考虑了浮点误差1e-3进行比较，保证并行卷积算法的正确性。

#include <CL/sycl.hpp>
#include <iostream>
#include <chrono>
using namespace sycl;

#define random_float() (rand() / double(RAND_MAX))
const int M = 10;
const int N = 10;
const int K = 3;

std::vector<float> Conv2D_Normal(const std::vector<float>& img, const std::vector<float>& kernel, double* dur) {
    auto s = std::chrono::high_resolution_clock::now();
    std::vector<float> result(M*N, 0.0f);

    for (int x = 0; x < M; ++x) {
        for (int y = 0; y < N; ++y) {
            float sum = 0.0f;
            for (int i = -K/2; i <= K/2; ++i) {
                for (int j = -K/2; j <= K/2; ++j) {
                    int xi = x + i;
                    int yj = y + j;
                    if (xi >= 0 && xi < M && yj >= 0 && yj < N) {
                        sum += img[xi*N+yj] * kernel[(i+K/2)*K+(j+K/2)];
                    }
                }
            }
            result[x*N+y] = sum;
        }
    }
    
    auto e = std::chrono::high_resolution_clock::now();
    auto duration = std::chrono::duration<float, std::milli>(e - s).count();
    if(dur) *dur+=duration;
    return result;
}


std::vector<float> Conv2D_Parallel(const std::vector<float>& img, const std::vector<float>& kernel, sycl::queue &q, double* dur) {
    std::vector<float> result(M*N, 0.0f);
    buffer img_buf(img.data(), range(M*N));
    buffer kernel_buf(kernel.data(), range(K*K));
    buffer result_buf(result.data(), range(M*N));
    
    auto e = q.submit([&](handler& h) {
        auto img_acc = img_buf.get_access<access::mode::read>(h);
        auto kernel_acc = kernel_buf.get_access<access::mode::read>(h);
        auto result_acc = result_buf.get_access<access::mode::write>(h);

        h.parallel_for(range(M, N), [=](id<2> idx) {
            int x = idx[0];
            int y = idx[1];
            float sum = 0.0f;

            for (int i = -K/2; i <= K/2; ++i) {
                for (int j = -K/2; j <= K/2; ++j) {
                    int xi = x + i;
                    int yj = y + j;
                    if (xi >= 0 && xi < M && yj >= 0 && yj < N) {
                        sum += img_acc[xi*N + yj] * kernel_acc[(i+K/2)*K + (j+K/2)];
                    }
                }
            }

            result_acc[x*N + y] = sum;
        });
    });
    e.wait();
    auto duration = (e.get_profiling_info<info::event_profiling::command_end>() -e.get_profiling_info<info::event_profiling::command_start>()) /1000.0f/1000.0f;
    if(dur) *dur+=duration;
    return result;
}

int main() {
    std::vector<float> img(M*N, 1.0f);
    std::vector<float> kernel(K*K, 1.0f);
    for(int i = 0; i < M*N; i++){
        img[i] = random_float();
    }
    for(int i = 0; i < K*K; i++){
        kernel[i] = random_float();
    }
    auto propList = sycl::property_list {sycl::property::queue::enable_profiling()};
    queue my_queue(default_selector_v , propList);
    
    // verify
    auto result1 = Conv2D_Normal(img,kernel,nullptr);
    auto result2 = Conv2D_Parallel(img,kernel,my_queue,nullptr);
    for(int i = 0; i < 10; i++){
        for(int j = 0; j < 10; j++){
            //std::cout<<result1[i*M+j]<<","<<result2[i*M+j]<<std::endl;
            if(abs(result1[i*M+j] - result2[i*M+j])>1e-3){
                std::cout<<"error"<<std::endl;
                return 0;
            }
        }
    }
    std::cout<<"the same"<<std::endl;
    
    // timer
    int iterations = 10;
    int warmup = 10;
    double duration_parallel = 0.0f;
    double duration_normal = 0.0f;
    for(int run = 0; run<iterations+warmup; run++){
        Conv2D_Parallel(img,kernel,my_queue,&duration_parallel);
    }
    duration_parallel/=iterations;
    
    for(int run = 0; run<iterations+warmup; run++){
        Conv2D_Normal(img,kernel,&duration_normal);
    }
    duration_normal/=iterations;
    
    std::cout<<"parallel: "<<duration_parallel<<" ms, normal: "<<duration_normal<<" ms.\n";
    

    return 0;
}

实验结果

执行编译脚本icpx -fsycl convolution.cpp -o convolution，然后运行convolution可执行文件，结果如下：

image.png

并行卷积的结果是对的，但是并行的耗时甚至远高于普通卷积，个人分析是由于在这个实验中我只是简单的将卷积并行化，没有做分块操作，而实验中使用的机器只有CPU core，也没有100个用于并行执行任务的线程，所以并行数量太高反而导致了竞争和通信调度的开销。由于本实验设置的M=10，N=10，所以100个双精度浮点数的存储空间应该不会超过缓存大小，计算过程中所有数据都能被装进data cache中，这个劣化的并行结果并不是图像太大导致频繁换入换出导致的。

总结

在本次实验中我学会了使用Intel OneAPI进行并行编程计算二维卷积。卷积运算是深度学习中非常重要的基本工具，卷积的每一位的结果之间互相不依赖（对比矩阵乘法），因此卷积运算有很大的并行化优化空间，且是最在各种异构的计算设备上（GPU/CPU/TPU/FPGA）运行的运算。

Intel OneAPI的设备选择器selector是一个新颖的地方，不同的selector匹配具有不同属性的计算设备，因此可以执行不同的计算方法。

通过本次实验，我使用OneAPI提供的设备计时器发现了在并行数过高，远远超过了计算设备的能力时，性能不升反降的现象，并对此进行了分析。OneAPI同样提供了分析工具例如Intel® oneAPI AI Analytics Toolkit，可以对代码的热点性能部分进行分析，提示用户将这些部分并行化加速。

综上，我认为Intel OneAPI所能提供的不仅仅是一套并行计算开发工具，还可以作为并行计算/深度学习的学习辅助入门工具，以及现有代码的性能检测优化工具。

avatar
文章
阅读
粉丝