oneAPI并行编程实践

上学日记

已于 2023-12-03 18:28:05 修改

阅读量821

点赞数 31

文章标签： oneapi c++

于 2023-12-03 17:16:31 首次发布

本文链接：https://blog.csdn.net/weixin_44340947/article/details/134766198

版权

文章目录

前言

在当今快速发展的计算世界中，异构计算不仅是一种趋势，更是实现高效、可扩展计算的必经之路。最近，我有幸深入学习了英特尔的oneAPI —— 一个旨在简化多平台软件开发的编程模型。在具体实践中，我选择了C++/SYCL作为我的主要编程工具，这是oneAPI的核心组成部分之一。SYCL是一个基于标准C++的高级编程模型，它为异构计算提供了强大的支持。通过使用SYCL，我们能够将复杂的并行计算任务直观且高效地映射到多种硬件上，从而充分利用了现代硬件的性能潜力。在本篇博客中，我想分享我在使用oneAPI和C++/SYCL解决实际问题时的心得体会。

作业环境

作业环境使用英特尔oneAPI Developer Cloud 服务,直接利用Developer Cloud平台中的CPU与GPU硬件完成相应的作业。

项目一：并⾏矩阵乘法

问题描述

编写⼀个基于oneAPI的C++/SYCL程序来执行矩阵乘法操作。需要考虑大尺寸矩阵的乘法操作以及不同线程之间的数据依赖关系。通常在实现矩阵乘法时，可以使用块矩阵乘法以及共享内存来提高计算效率。

对已有代码的分析

在 JupyterLab 生成的默认文件列表中有一个SYCL_Performance_Portability 文件夹，实现了基于SYCL的GEMM，并对比了算法在不同平台上的性能表现，主要代码如下图所示，代码结构包括两部分：Kernel Code 和 Common Code
在这里插入图片描述

Common Code

通用代码即图中的 /lab/mm_dpcpp_common.cpp 包含了命令行参数解析、初始化矩阵、设置队列以及调用核心代码、结果验证等等，代码如下：

//==============================================================
// Matrix Multiplication: SYCL Matrix Multiplication Common
//==============================================================
// Copyright © 2021 Intel Corporation
//
// SPDX-License-Identifier: MIT
// =============================================================


#include <CL/sycl.hpp>
#include <getopt.h>
#include <ctime>
#include <chrono>

using namespace sycl;

//# matrix multiplication kernel implementation in mm_dpcpp_*.cpp
void mm_kernel(queue &q, std::vector<float> &matrix_a, std::vector<float> &matrix_b, std::vector<float> &matrix_c, size_t N, size_t M);

//# floating point error verification function
bool almost_equal(float a, float b){
    float tolerance = 1e-6;
    float diff = fabs(a - b);
    a = fabs(a);
    b = fabs(b);
    float bigger = (b > a) ? b : a;
    if(diff <= bigger * tolerance) return true;
    return false;
}

int main(int argc, char *argv[]) {
    
    size_t N = 1024;
    size_t M = 16;
    int VERIFY = 0;
    int PRINT_OUTPUT_MATRIX = 0;

    //# command line arguments
    int arg;
    while ((arg = getopt (argc, argv, "n:m:vp")) != -1)
        switch (arg){
            case 'n':
                N = std::atoi(optarg);
                break;
            case 'm':
                M = std::atoi(optarg);
                break;
            case 'v':
                VERIFY = 1;
                break;
            case 'p':
                PRINT_OUTPUT_MATRIX = 1;
                break;
            case 'h':
                std::cout << std::endl;
                std::cout << "Usage   : ./a.out -n <MATRIX_SIZE> -m <WORK_GROUP_SIZE> -v -p\n\n";
                std::cout << "          [-n] size for matrix, eg: 1024\n";
                std::cout << "          [-m] size of work_group, eg: 8/16\n";
                std::cout << "          [-v] verify output with linear computation on cpu\n";
                std::cout << "          [-p] print output matrix\n";
                std::cout << "Example : ./a.out -n 1024 -m 16 -v -p\n\n";
                std::exit(0);
        }

    //# Define vectors for matrices
    std::vector<float> matrix_a(N*N);
    std::vector<float> matrix_b(N*N);
    std::vector<float> matrix_c(N*N);
    std::vector<float> matrix_d(N*N);
    
    //# Initialize matrices with values
    float v1 = 2.f;
    float v2 = 3.f;
    for (int i=0; i<N; i++)
        for (int j=0; j<N; j++){
            matrix_a[i*N+j] = v1++;
            matrix_b[i*N+j] = v2++;
            matrix_c[i*N+j] = 0.f;
            matrix_d[i*N+j] = 0.f;
    }
    
    //# Define queue with default device for offloading computation
    queue q(property::queue::enable_profiling{});
    std::cout << "Offload Device        : " << q.get_device().get_info<info::device::name>() << "\n";
    std::cout << "max_work_group_size   : " << q.get_device().get_info<info::device::max_work_group_size>() << "\n";
    
    //# get start time
    auto start = std::chrono::high_resolution_clock::now().time_since_epoch().count();
    
    //# Call matrix multiplication kernel implementation
    mm_kernel(q, matrix_a, matrix_b, matrix_c, N, M);
    
    //# print kernel compute duration from host
    auto duration = std::chrono::high_resolution_clock::now().time_since_epoch().count() - start;
    std::cout << "Compute Duration      : " << duration / 1e+9 << " seconds\n";
    
    //# Print Output if -p in cmd-line
    if (PRINT_OUTPUT_MATRIX){
        for (int i=0; i<N; i++){
            for (int j=0; j<N; j++){
                std::cout << matrix_c[i*N+j] << " ";
            }
            std::cout << "\n";
        }
    } else {
        std::cout << " [0][0] = " << matrix_c[0] << "\n";
    }
    
    //# Compute local and compare with offload computation if -v in cmd-line
    if (VERIFY){
        int fail = 0;
        for(int i=0; i<N; i++){
            for (int j = 0; j < N; j++) {
                for(int k=0; k<N; k++){
                    matrix_d[i*N+j] += matrix_a[i*N+k] * matrix_b[k*N+j];
                }
                if(!almost_equal(matrix_c[i*N+j], matrix_d[i*N+j])) fail = 1;
            }
        }
        if(fail == 1){
            std::cout << "FAIL\n";
        } else {
            std::cout << "PASS\n";
        }
    }
    return 0;
}

Kernel Code

矩阵乘法计算的核心代码放在一个单独的源文件中，该文件由通用代码源调用。lab文件夹中包含了多个不同方法实现并行计算的核心代码，主要分析如下：

Basic Version (mm_dpcpp_basic.cpp):
- 使用了一个基本的并行模式，没有指定工作组大小。
- 在mm_kernel函数中，通过parallel_for(range<2>{N,N}, [=](item<2> item){...})来设置并行执行的范围，其中range<2>{N,N}指定了全局尺寸。
NDRange (mm_dpcpp_ndrange.cpp):
- 在mm_kernel函数中，parallel_for(nd_range<2>{global_size, work_group_size}, [=](nd_item<2> item){...})来设置并行执行的范围，其中nd_range<2>{global_size, work_group_size}定义了全局尺寸和工作组尺寸。
- 通过显式指定工作组的大小，为开发者提供了更多的控制和优化空间。
NDRange Variable (mm_dpcpp_ndrange_var.cpp):
- 类似于NDRange实现，但在计算过程中引入了局部变量来存储中间结果。
- 在parallel_for内部，使用临时变量temp来累加乘法结果，最后将temp的值赋给输出矩阵C的对应元素。
- 这种方法优化了对全局内存的访问，减少了内存写操作的次数。
Local Memory (mm_dpcpp_localmem.cpp):
- 使用了SYCL的本地内存访问器（local accessors）。
- 在mm_kernel函数中，除了全局和工作组尺寸的设置，还额外定义了本地内存访问器A_tile和B_tile，用于存储工作组内的局部矩阵数据。
- 这种方法利用了本地内存来提高数据访问的效率。
MKL (mm_dpcpp_mkl.cpp):
- 使用oneAPI Math Kernel Library (oneMKL)实现矩阵乘法。
- mm_kernel函数中，创建了矩阵的缓冲区并调用oneapi::mkl::blas::gemm函数执行矩阵乘法。这是一个高级库调用，意味着性能优化和硬件加速由oneMKL处理。

使用USM实现矩阵乘法

统一共享内存（Unified Shared Memory, USM）在SYCL中提供了一种不同于传统缓冲区（buffer）或访问器（accessor）方法的内存管理方式。使用USM的优点在于简化了内存管理，因为它允许CPU和GPU访问相同的内存地址，减少了数据复制和同步的需要。

在~/SYCL_Performance_Portability/lab 中新建mm_dpcpp_usm.cpp ，代码如下，将Kernel Code 和 Common Code写在了一个源文件

//==============================================================
// Matrix Multiplication: SYCL USM
//==============================================================
// Copyright © 2021 Intel Corporation
//
// SPDX-License-Identifier: MIT
// =============================================================

#include <CL/sycl.hpp>
#include <getopt.h>
#include <ctime>
#include <chrono>

using namespace sycl;

//# matrix multiplication kernel implementation in mm_dpcpp_*.cpp
void mm_kernel(queue &q, float *a, float *b, float *c, size_t N, size_t M) {
    // 提交任务到队列
		event e = q.submit([&](handler &h) {
        //# Define size for ND-Range and work-group size
        range<2> global_size(N,N);
        range<2> work_group_size(M,M);
        h.parallel_for(nd_range<2>{global_size, work_group_size}, [=](nd_item<2> item){
            const int i = item.get_global_id(0);
            const int j = item.get_global_id(1);
            //# Use private mem to store intermediate result
            float temp = 0.f;
            for (int k = 0; k < N; k++) {
                temp += a[i*N+k] * b[k*N+j];
            }
            c[i*N+j] = temp;
        });
    });
    e.wait(); // 等待队列中的任务完成
    auto kernel_duration = (e.get_profiling_info<info::event_profiling::command_end>() -
                            e.get_profiling_info<info::event_profiling::command_start>());
    std::cout << "Kernel Execution Time : " << kernel_duration / 1e+9 << " seconds\n";
}

//# floating point error verification function
bool almost_equal(float a, float b){
    float tolerance = 1e-6;
    float diff = fabs(a - b);
    a = fabs(a);
    b = fabs(b);
    float bigger = (b > a) ? b : a;
    if(diff <= bigger * tolerance) return true;
    return false;
}

int main(int argc, char *argv[]) {
    
    size_t N = 1024;
    size_t M = 16;
    int VERIFY = 0;
    int PRINT_OUTPUT_MATRIX = 0;

    //# command line arguments
    int arg;
    while ((arg = getopt (argc, argv, "n:m:vp")) != -1)
        switch (arg){
            case 'n':
                N = std::atoi(optarg);
                break;
            case 'm':
                M = std::atoi(optarg);
                break;
            case 'v':
                VERIFY = 1;
                break;
            case 'p':
                PRINT_OUTPUT_MATRIX = 1;
                break;
            case 'h':
                std::cout << std::endl;
                std::cout << "Usage   : ./a.out -n <MATRIX_SIZE> -m <WORK_GROUP_SIZE> -v -p\\n\\n";
                std::cout << "          [-n] size for matrix, eg: 1024\\n";
                std::cout << "          [-m] size of work_group, eg: 8/16\\n";
                std::cout << "          [-v] verify output with linear computation on cpu\\n";
                std::cout << "          [-p] print output matrix\\n";
                std::cout << "Example : ./a.out -n 1024 -m 16 -v -p\\n";
                return 0;
        }
    //# Define queue with default device for offloading computation
    queue q(property::queue::enable_profiling{});
    
    //# Allocate memory using USM
    float *matrix_a = malloc_shared<float>(N * N, q);
    float *matrix_b = malloc_shared<float>(N * N, q);
    float *matrix_c = malloc_shared<float>(N * N, q);
    float *matrix_d = nullptr; // Only allocate if verification is needed

    if (VERIFY) {
        matrix_d = malloc_shared<float>(N * N, q);
    }

    // Initialize matrices with values
    float v1 = 2.f;
    float v2 = 3.f;
    for (int i=0; i<N; i++)
        for (int j=0; j<N; j++){
            matrix_a[i*N+j] = v1++;
            matrix_b[i*N+j] = v2++;
            matrix_c[i*N+j] = 0.f;
            if (VERIFY) {
                matrix_d[i*N+j] = 0.f;
            }
        }


    std::cout << "Offload Device        : " << q.get_device().get_info<info::device::name>() << "\n";
    std::cout << "max_work_group_size   : " << q.get_device().get_info<info::device::max_work_group_size>() << "\n";
    //# get start time
    auto start = std::chrono::high_resolution_clock::now().time_since_epoch().count();
    
    //# Call matrix multiplication kernel implementation
    mm_kernel(q, matrix_a, matrix_b, matrix_c, N, M);

    //# print kernel compute duration from host
    auto duration = std::chrono::high_resolution_clock::now().time_since_epoch().count() - start;
    std::cout << "Compute Duration      : " << duration / 1e+9 << " seconds\n";
    
    //# Print Output if -p in cmd-line
    if (PRINT_OUTPUT_MATRIX){
        for (int i=0; i<N; i++) {
            for (int j=0; j<N; j++) {
                std::cout << matrix_c[i*N+j] << " ";
            }
            std::cout << "\\n";
        }
    } else {
        std::cout << " [0][0] = " << matrix_c[0] << "\\n";
    }

    //# Compute local and compare with offload computation if -v in cmd-line
    if (VERIFY){
        int fail = 0;
        for(int i=0; i<N; i++){
            for (int j = 0; j < N; j++) {
                for(int k=0; k<N; k++){
                    matrix_d[i*N+j] += matrix_a[i*N+k] * matrix_b[k*N+j];
                }
                if(!almost_equal(matrix_c[i*N+j], matrix_d[i*N+j])) fail = 1;
            }
        }
        if(fail == 1){
            std::cout << "FAIL\\n";
        } else {
            std::cout << "PASS\n";
        }
    }
    
    free(matrix_a, q);
    free(matrix_b, q);
    free(matrix_c, q);
    if (VERIFY) {
        free(matrix_d, q);
    }
    return 0;
}

运行结果

在~/SYCL_Performance_Portability/run_all.sh文件末中加入：

echo ====================
echo mm_dpcpp_usm
dpcpp ${src}mm_dpcpp_usm.cpp -o ${src}mm_dpcpp_usm -w -O3
./${src}mm_dpcpp_usm$arg

打开终端，在~/SYCL_Performance_Portability 目录下执行命令./q run_all.sh "GPU GEN9" ，将Intel® Gen9 GPU作为device，如下图

在这里插入图片描述

得到结果如下

在这里插入图片描述

我们看到在该环境下，MATRIX_SIZE= 1024x1024 ， WORK_GROUP_SIZE= 16x16 时，几种方法的性能排行是：

Local Memory > USM > NDRange Variable > NDRange > MKL

项目二：并⾏排序算法

问题描述

使用基于oneAPI的C++/SYCL实现⼀个高效的并行归并排序。需要考虑数据的分割和合并以及线程之间的协作。

使用USM实现并行排序

在项目二中，我们借鉴项目一使用USM分配用于存储数据的内存，实现了并行归并排序，并验证了排序结果的正确性，核心设计如下：

分解任务:
- 并行归并排序通过将大数组分解为较小的子数组来实现并行化。这些子数组可以并行排序。
使用nd_range:
- 在SYCL中，使用nd_range对象来定义全局和局部工作项的数量。这允许您精确控制工作分配，以便优化对特定硬件的性能。
- 全局大小（globalSize）定义了总的工作项数量，而局部大小（localSize）定义了每个工作组中的工作项数量。
并行执行:
- 在parallelMergeSort函数中，通过q.submit将归并排序的任务提交到SYCL队列。
- 使用parallel_for结合nd_range，将任务分配给多个工作项。每个工作项负责数组的一个部分。
合并操作:
- 每次迭代的合并操作是并行执行的。合并操作将两个已排序的子数组合并为一个有序数组，这一过程在多个工作项之间分配。
数据同步:
- 在每次合并操作之后，需要确保所有工作项都完成了它们的任务，以便可以进行下一轮合并。在SYCL中，这通过在每次提交任务后调用q.wait()来实现。

sort.cpp完整代码如下：

#include <CL/sycl.hpp>
#include <iostream>
#include <vector>
#include <algorithm>
#include <chrono>

using namespace sycl;

void parallelMergeSort(queue &q, float* data, size_t N, size_t M) {
    float* temp = malloc_shared<float>(N, q);

    const size_t localSize = M; 
    const size_t globalSize = ((N + localSize - 1) / localSize) * localSize;

    for (size_t width = 1; width < N; width *= 2) {
        q.submit([&](handler &h) {
            h.parallel_for(nd_range<1>(range<1>(globalSize), range<1>(localSize)), [=](nd_item<1> item) {
                size_t idx = item.get_global_id(0);
                size_t left = 2 * width * idx;
                size_t middle = std::min(left + width, N);
                size_t right = std::min(left + 2 * width, N);

                if (left < N) {
                    size_t i = left, j = middle, k = left;

                    while (i < middle && j < right) {
                        if (data[i] < data[j]) {
                            temp[k++] = data[i++];
                        } else {
                            temp[k++] = data[j++];
                        }
                    }

                    while (i < middle) {
                        temp[k++] = data[i++];
                    }

                    while (j < right) {
                        temp[k++] = data[j++];
                    }
                }
            });
        });

        // 将临时数组中的数据复制回原数组
        q.submit([&](handler &h) {
            h.parallel_for(range<1>(N), [=](id<1> idx) {
                data[idx] = temp[idx];
            });
        });
    }

    q.wait();

    // 释放临时数组内存
    free(temp, q);
}


int main() {
    size_t N = 65536; // 数据大小
    size_t M = 16;  // 每个工作组的工作项数
    queue q;

    // 使用USM分配共享内存
    float* data = malloc_shared<float>(N, q);
    std::vector<float> serialData(N); // 用于串行排序的数据副本

    // 初始化数组数据
    for (size_t i = 0; i < N; ++i) {
        float val = rand() % 1000; // 生成0到999之间的随机数
        data[i] = val;
        serialData[i] = val;
    }

    // 并行排序时间测量
    auto startParallel = std::chrono::high_resolution_clock::now().time_since_epoch().count();
    parallelMergeSort(q, data, N, M);
    auto endParallel = std::chrono::high_resolution_clock::now().time_since_epoch().count();
    double durationParallel = (endParallel - startParallel) / 1e+9; // 转换为秒
    
    std::cout << "Offload Device        : " << q.get_device().get_info<info::device::name>() << "\n";
    std::cout << "max_work_group_size   : " << q.get_device().get_info<info::device::max_work_group_size>() << "\n";
    
    // 输出结果
    std::cout << "并行排序时间 : " << durationParallel << " 秒.\n";

    // 验证排序结果
    std::sort(serialData.begin(), serialData.end());
    bool isCorrect = true;
    for (size_t i = 0; i < N; ++i) {
        if (data[i] != serialData[i]) {
            isCorrect = false;
            break;
        }
    }

    if (isCorrect) {
        std::cout << "验证结果: 排序正确\n";
    } else {
        std::cout << "验证结果: 排序错误\n";
    }

    free(data, q);
    return 0;
}

运行结果

在~/SYCL_Performance_Portability/ 创建文件run_sort.sh：

#!/bin/bash
source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
src="lab/"

echo ====================
echo sort
icpx -fsycl ${src}sort.cpp -o ${src}sort -w -O3
./${src}sort

执行命令./q run_sort.sh "GPU GEN9" ，将Intel® Gen9 GPU作为device，在随机生成的 N = 65536 的数据大小情况下，运行结果如下图:

在这里插入图片描述

项目三：图像卷积并⾏加速

描述

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

分析

图像卷积是一种常见的图像处理操作，用于应用各种滤波器和特征检测器。其原理可以简单地描述为在图像的每个像素上应用一个小的矩阵（通常称为卷积核或滤波器），并将卷积核中的元素与图像中对应位置的像素值相乘，然后将所有乘积的和作为结果。这个过程可以看作是对图像进行了平滑、锐化、边缘检测等操作。
假设有⼀个大小为M × N 的输入图像I 和一个大小为m × n 的卷积核 K 。图像卷积操作可以用下面的数学公式来表示：

$\ S(i, j) = \sum_{k=0}^{m-1} \sum_{l=0}^{n-1} I(i + k, j + l) \cdot K(k, l) \$
这里：

$S (i, j)$ 是结果图像中位置 (i, j) 的像素值。
$I (i + k, j + l)$ 是输入图像中位置 (i + k, j + l) 的像素值。
$K (k, l)$ 是卷积核中位置 (k, l) 的权重。
在这个过程中，卷积核 K 在输入图像 I 上逐个像素地滑动，对于每个位置 (i, j)，卷积核覆盖的 I 的区域与 K 对应位置的元素相乘，然后将所有这些乘积相加以获得 S 的相应位置的值。

卷积核通常是一个小的⼆维矩阵，用于捕捉图像中的特定特征。在异构计算编程中，可以使用并行计算来加速图像卷积操作。通过将图像分割成小块，然后在GPU上并行处理这些块，可以实现高效的图像卷积计算。通过合理的块大小和线程组织方式，可以最大限度地利用GPU的并行计算能力来加速图像处理过程。

使用USM实现图像卷积操作

基于GPU的图像卷积操作与矩阵乘法类似，设计如下：

图像卷积核函数 convolve_kernel：
- 接收图像、卷积核、输出数组和它们的尺寸。
- 使用 SYCL 的统一共享内存（USM）分配和管理内存。
- 使用 nd_range 来定义工作组和工作项的布局，优化 GPU 上的并行性能。
- 在 GPU 上执行实际的卷积操作，每个工作项处理图像的一部分。
主函数 main：
- 定义图像、卷积核的大小和数据。
- 使用随机数填充图像和卷积核以模拟实际情况。
- 创建一个 SYCL 队列，用于将计算任务提交给 GPU。
- 调用 convolve_kernel 函数执行卷积，并测量执行时间。
性能测量：
- 使用 std::chrono 库在卷积操作前后记录时间，计算并输出执行所需的时间。

conv.cpp 代码如下：

#include <CL/sycl.hpp>
#include <vector>
#include <iostream>
#include <random>
#include <chrono>

using namespace sycl;

// 图像卷积核函数
void convolve_kernel(queue &q, float *image, float *kernel, float *output, size_t M, size_t N, size_t m, size_t n) {
    // 使用 USM 分配内存
    auto image_d = malloc_shared<float>(M * N, q);
    auto kernel_d = malloc_shared<float>(m * n, q);
    auto output_d = malloc_shared<float>(M * N, q);

    // 拷贝数据到设备内存
    std::memcpy(image_d, image, sizeof(float) * M * N);
    std::memcpy(kernel_d, kernel, sizeof(float) * m * n);
    
     // 定义 nd_range
    const size_t local_size = 16; // 本地工作组大小
    range<2> local_range(local_size, local_size);
    range<2> global_range((M + local_size - 1) / local_size * local_size, (N + local_size - 1) / local_size * local_size);
    nd_range<2> ndRange(global_range, local_range);

    // 记录开始时间
    auto start = std::chrono::high_resolution_clock::now();

    // 提交任务到队列
    q.parallel_for(ndRange, [=](nd_item<2> item) {
        size_t i = item.get_global_id(0);
        size_t j = item.get_global_id(1);
        if (i < M && j < N) {
            float sum = 0.0f;
            // 卷积操作
            for (size_t ki = 0; ki < m; ++ki) {
                for (size_t kj = 0; kj < n; ++kj) {
                    size_t ii = i + ki - m / 2;
                    size_t jj = j + kj - n / 2;
                    if (ii >= 0 && ii < M && jj >= 0 && jj < N) {
                        sum += image_d[ii * N + jj] * kernel_d[ki * n + kj];
                    }
                }
            }
            output_d[i * N + j] = sum;
        }
    }).wait();


    std::cout << "Offload Device        : " << q.get_device().get_info<info::device::name>() << "\n";
    std::cout << "max_work_group_size   : " << q.get_device().get_info<info::device::max_work_group_size>() << "\n";
    // 记录结束时间
    auto end = std::chrono::high_resolution_clock::now();

    // 计算并输出执行时间
    auto duration = duration_cast<std::chrono::microseconds>(end - start);
    std::cout << "Compute Duration      :" << duration.count() / 1e+6 << " s" << std::endl;

    // 将结果拷贝回主机
    std::memcpy(output, output_d, sizeof(float) * M * N);

    // 释放设备内存
    free(image_d, q);
    free(kernel_d, q);
    free(output_d, q);
}

int main() {
    // 图像和卷积核的大小
    size_t M = 1024, N = 1024, m = 3, n = 3;

    // 初始化数据
    std::vector<float> image(M * N);
    std::vector<float> kernel(m * n);
    std::vector<float> output(M * N, 0.0f);

    // 随机数生成器
    std::mt19937 rng;
    rng.seed(std::random_device()());
    std::uniform_real_distribution <float> dist(0.0, 1.0);

    // 使用随机数填充图像和卷积核
    for (auto& val : image) {
        val = dist(rng);
    }
    for (auto& val : kernel) {
        val = dist(rng);
    }

    // 创建队列
    queue q;

    // 执行卷积核函数
    convolve_kernel(q, image.data(), kernel.data(), output.data(), M, N, m, n);

    return 0;
}

运行结果

在~/SYCL_Performance_Portability/ 创建文件run_conv.sh：

#!/bin/bash
source /opt/intel/inteloneapi/setvars.sh > /dev/null 2>&1
src="lab/"

echo ====================
echo conv
icpx -fsycl ${src}conv.cpp -o ${src}conv -w -O3
./${src}conv

执行命令./q run_conv.sh "GPU GEN9" ，将Intel® Gen9 GPU作为device，在M = 1024, N = 1024, m = 3, n = 3 并随机生成数据的前提下，结果如下图:

在这里插入图片描述

总结

通过深入研究和实践oneAPI在各种并行计算场景中的应用，我获得了宝贵的知识和经验。 oneAPI作为一个先进的编程模型，不仅增强了我的对并行计算概念的理解，也让我认识到了跨平台编程在现代计算中的重要性。通过这几个具体的作业案例，我学会了如何有效地利用SYCL和局部内存、nd_range等概念来优化性能，这对于未来面临的任何高性能计算挑战都是宝贵的经验。

oneAPI提供一致性编程接口，使得应用跨平台复用，对于解决当今面临的各种计算挑战至关重要。oneAPI不仅提供了强大的工具集合来应对这些挑战，还提供了一个统一的平台来简化复杂的编程任务，从而使开发者能够更专注于解决实际问题而非底层硬件的细节。

在这里插入图片描述

本次学习不仅仅是学习一个新工具的过程，更是一次深入理解并行计算原理和实践的旅程。随着计算需求的不断增长和技术的不断进步，我相信，像oneAPI这样的先进工具将继续在推动科技创新的前沿发挥关键作用。

上学日记

关注

31
点赞
踩
31

收藏

觉得还不错? 一键收藏
2
评论
oneAPI并行编程实践

在当今快速发展的计算世界中，异构计算不仅是一种趋势，更是实现高效、可扩展计算的必经之路。最近，我有幸深入学习了英特尔的oneAPI —— 一个旨在简化多平台软件开发的编程模型。在具体实践中，我选择了C++/SYCL作为我的主要编程工具，这是oneAPI的核心组成部分之一。SYCL是一个基于标准C++的高级编程模型，它为异构计算提供了强大的支持。通过使用SYCL，我们能够将复杂的并行计算任务直观且高效地映射到多种硬件上，从而充分利用了现代硬件的性能潜力。
复制链接

扫一扫