团队名称：中南ovo
一、引言
本次实验报告源于中南大学intel校企合作，我在C++课上对intel oneAPI jupyterlab 进行了学习，在这上面了解了SYCL相关知识，使用英特尔Developer Cloud平台进行矩阵乘法的异构计算，并探索不同统一共享内存（USM）类型对性能的影响。
进行本次实验需要在C++的基础上，了解SYCL编程知识，建议阅读前先进行SYCL编程学习。
二、oneAPI平台简单介绍
oneAPI 是一个由英特尔（Intel）推出的跨平台编程模型，旨在简化和统一多种计算架构（如CPU、GPU、FPGA等）的编程。其核心目标是提供一种能够在多种硬件上运行的统一编程方法，以减少开发者在针对不同硬件编写和优化代码时的工作量。
oneAPI 包括一系列的工具和库，如：
- DPC++（Data Parallel C++）：一种基于C++的编程语言，用于编写可以在多种硬件架构上运行的并行代码。
- Level Zero：一种低级别的硬件抽象层，用于直接与硬件交互。
- oneDNN、oneTBB等库：提供针对深度学习、线程构建块等的优化。
通过这些工具和库，oneAPI 旨在提高开发效率，同时保持高性能的计算能力，使开发者能够更容易。
三、实验
（1）环境配置：推荐使用英特尔oneAPI Developer Cloud 服务，可免安装额外环境，直接利用Developer Cloud平台中的CPU与GPU硬件完成相应的作业。 相应注册及启用服务可参考相应的使用手册，关于在登录节点上运行的JupyterLab递交代码到计算节点的排队系统的相关操作方式及命令，可以参考任务递交相应命令参考 以及 队列管理参考 。
如果使用自有硬件，推荐使用基于英特尔酷睿6代或更新版本的处理器与系统中同时可用的基于Xe或Arc的集成显卡，通过完成与安装操作系统配套的相应oneAPI软件工具，通过构建自有开发环境方式完成作业。
相应的工具可参考oneAPI基础工具套件或独立编译器的网页， 硬件要求可参考发布说明中的具体要求。
（2）在开始进行实验之前，建议先阅读intel oneAPI jupyterlab 平台“SYCL_PROGRAM_STRUCTURE”文档进行学习，网址：JupyterHub
我们先来看文档中的 vector add实验
Lab Exercise: Vector Add
Complete the coding excercise below using SYCL Buffer and Accessor concepts:
- The code has three vector
vector1initialized on host - The kernel code increments the
vector1by 1. - Create a new second
vector2and initialize to value 20. - Create sycl buffers for the above second vector
- In the kernel code, create a second accessor for the second vector buffer
- Modify the vector increment to vector add, by adding
vector2tovector1
- Edit the code cell below by following the steps and then click run ▶ to save the code to a file.
- Next run ▶ the cell in the Build and Run section below the code to compile and execute the code.
#include <sycl/sycl.hpp>
using namespace sycl;
int main() {
const int N = 256;
//# Initialize a vector and print values
std::vector<int> vector1(N, 10);
std::cout << "\nInput Vector1: ";
for (int i = 0; i < N; i++) std::cout << vector1[i] << " ";
//# STEP 1 : Create second vector, initialize to 20 and print values
std::vector<int> vector2(N, 20);
std::cout << "\nInput Vector2: ";
for (int i = 0; i < N; i++) std::cout << vector2[i] << " ";
//# Create Buffer
buffer vector1_buffer(vector1);
//# STEP 2 : Create buffer for second vector
buffer vector2_buffer(vector2);
//# Submit task to add vector
queue q;
q.submit([&](handler &h) {
//# Create accessor for vector1_buffer
accessor vector1_accessor(vector1_buffer, h);
//# STEP 3 - add second accessor for second buffer
accessor vector2_accessor(vector2_buffer, h, read_only);
h.parallel_for(range<1>(N), [=](id<1> index) {
//# STEP 4 : Modify the code below to add the second vector to first one
vector1_accessor[index] += vector2_accessor[index];
});
});
//# Create a host accessor to copy data from device to host
host_accessor h_a(vector1_buffer, read_only);
//# Print Output values
std::cout << "\nOutput Values: ";
for (int i = 0; i < N; i++) std::cout << h_a[i] << " ";
std::cout << "\n";
return 0;
}
这个小实验是我们进行sycl编程的开始，它的核心是用sycl编写两个向量相加，并在异构计算环境中运行这个程序。
在刚接触sycl的情况下，我认为在实验中主要有以下难点：
1.理解并行计算的概念
在SYCL中，kernel是在device（如GPU）上并行执行的。这就要求我们思考，如何充分利用device的计算资源，例如，将原来C++程序中的循环结构改成device中并行计算的kernel。
2.熟悉SYCL的缓冲区和访问器
-
缓冲区（Buffer）：缓冲区是SYCL中用于存储数据并在主机和设备之间传输数据的对象。它抽象了内存管理的细节，允许SYCL运行时负责在主机和设备间透明地复制数据。
-
访问器（Accessor）：访问器是用于访问缓冲区中数据的对象。它提供了一种方法来读取或写入缓冲区中的数据，并可以指定访问类型（读、写或读写）和访问范围。
这里给出代码实例方便理解：// 创建缓冲区 buffer<float, 2> bufferA(A.data(), range<2>(N, N)); // 在内核函数中创建访问器 q.submit([&](handler& h) { auto accA = bufferA.get_access<access::mode::read>(h); auto accB = bufferB.get_access<access::mode::read>(h); auto accC = bufferC.get_access<access::mode::write>(h); // 使用访问器进行计算 h.parallel_for(range<2>(N, N), [=](id<2> id) { size_t i = id[0]; size_t j = id[1]; accC[i][j] = accA[i][j] + accB[i][j]; }); });在这个例子中，首先为矩阵A、B和C创建了二维缓冲区。然后，在内核函数中创建了这些缓冲区的访问器。这些访问器用于读取A和B的内容，以及写入C的结果。在内核中使用这些访问器来执行实际的向量加法操作。
理解了上面这两个难点，在熟悉相应的语法后，可以将add vector 实验补充完整：
#include <sycl/sycl.hpp>
using namespace sycl;
int main() {
const int N = 256;
//# Initialize a vector and print values
std::vector<int> vector1(N, 10);
std::cout << "\nInput Vector1: ";
for (int i = 0; i < N; i++) std::cout << vector1[i] << " ";
//# STEP 1 : Create second vector, initialize to 20 and print values
std::vector<int> vector2(N, 20);
std::cout << "\nInput Vector2: ";
for (int i = 0; i < N; i++) std::cout << vector2[i] << " ";
//# Create Buffer
buffer vector1_buffer(vector1);
//# STEP 2 : Create buffer for second vector
buffer vector2_buffer(vector2);
//# Submit task to add vector
queue q;
q.submit([&](handler &h) {
//# Create accessor for vector1_buffer
accessor vector1_accessor(vector1_buffer, h);
//# STEP 3 - add second accessor for second buffer
accessor vector2_accessor(vector2_buffer, h, read_only);
h.parallel_for(range<1>(N), [=](id<1> index) {
//# STEP 4 : Modify the code below to add the second vector to first one
vector1_accessor[index] += vector2_accessor[index];
});
});
//# Create a host accessor to copy data from device to host
host_accessor h_a(vector1_buffer, read_only);
//# Print Output values
std::cout << "\nOutput Values: ";
for (int i = 0; i < N; i++) std::cout << h_a[i] << " ";
std::cout << "\n";
return 0;
}
下图为在jupyterlab平台上运行的结果，可知结果正确：
这个小实验帮助我们理解了sycl中并行计算的概念，熟悉了sycl语法，下面让我们正式开始实验任务
（3）实验任务二：
矩阵乘法运算
按以下步骤写一个SYCL程序来实现包含两个矩阵的矩阵乘法运算：
组建2个序列（向量）的浮点数，每个序列的规格是N（如N=1024*1024），构成矩阵输入值。用随机值初始化序列。使用缓存和存储来实现对设备（GPU）上矩阵内存的分配并运行。运行SYCL Kernel实现两个矩阵的并行运算，这里你需要运用SYCL nd_range概念来定义Kernel的运行范围。使用SYCL排队系统来运行设备上的Kernel。 Kernel运行结束后，使用存储将结果从设备（GPU）检索回主机。
把这个任务分解成小任务，主要要攻克的点有以下部分：
1. 矩阵初始化
2. 缓冲区和访问器的使用
3. 编写SYCL Kernel进行矩阵乘法
4. 使用nd_range定义Kernel的运行范围
5. 使用SYCL排队系统和内核执行
6. 数据回传
接下来我们逐一分析这些关键点：
1.矩阵初始化
这里我们使用 <random> 库生成随机数并填充矩阵；并进行相应措施确保程序能够处理大量数据
void initialize_matrix(vector<float>& matrix) {
random_device rd;
mt19937 mersenne_engine(rd());
uniform_real_distribution<float> dist(-100.0, 100.0);
for (auto& element : matrix) {
element = dist(mersenne_engine);
}
}
2. 缓冲区和访问器的使用
这里的难点在于创建和管理SYCL缓冲区，正确使用访问器来读写缓冲区中的数据。
关键代码：
buffer<float, 2> bufferA(A.data(), range<2>(N, N));
这段代码创建了一个二维缓冲区 bufferA，用于存储矩阵A的数据。类似地，我们也可以按照相同方式为矩阵B和C创建缓冲区。这些缓冲区将在SYCL设备上使用。
3. 编写SYCL Kernel进行矩阵乘法
关键点：（1）处理索引和边界条件 （2）利用并行计算优势，分解矩阵乘法为可并行执行的子任务
关键代码：
h.parallel_for(work_range, [=](nd_item<2> item) {
size_t i = item.get_global_id(0);
size_t j = item.get_global_id(1);
float sum = 0.0f;
for (size_t k = 0; k < N; ++k) {
sum += accA[i * N + k] * accB[k * N + j];
}
accC[i * N + j] = sum;
});
4. 使用nd_range定义Kernel的运行范围
难点：理解nd_range概念，正确设置工作组大小和形状
关键代码：
nd_range<2> work_range(range<2>(N, N), range<2>(16, 16));
鉴于nd_range理解较难，这里我们具体讲讲每个参数含义：
-
二维索引空间：
nd_range<2>表示这是一个二维的工作索引空间。这里的“2”代表了工作项的维度，适用于像矩阵乘法这样的二维数据操作。 -
全局范围：
range<2>(N, N)定义了全局索引空间的大小，即在每个维度上有多少工作项。在这个例子中，每个维度的大小都被设置为了N，这对应于矩阵的行和列的大小。这意味着相应的kernel将会有N x N个工作项。 -
局部范围：
range<2>(16, 16)定义了每个工作组中的工作项的布局和数量。每个工作组包含16 x 16个工作项。
5. 使用SYCL排队系统和内核执行
难点：管理队列，确保Kernel按预期执行；处理执行和同步问题。
关键代码：
queue q{ property::queue::enable_profiling() };
auto event = q.submit([&](handler& h) { /* ... */ });
event.wait();
6. 数据回传
关键代码：
host_accessor h_accC(bufferC, read_only);
for (size_t i = 0; i < N; ++i) {
for (size_t j = 0; j < N; ++j) {
C[i * N + j] = h_accC[i][j];
}
}
这段代码中，使用了一个主机访问器 h_accC 来读取缓冲区 bufferC 的内容。这个访问器允许在主机端访问设备端的计算结果。然后通过一个双重循环，将计算结果从SYCL缓冲区复制回标准向量 C。
接下来给出完整的解决方案：
#include <CL/sycl.hpp>
#include <iostream>
#include <random>
#include <vector>
#include <chrono>
using namespace sycl;
using namespace std;
const size_t N = 1024;
// 用随机浮点数初始化矩阵的函数
void initialize_matrix(vector<float>& matrix) {
random_device rd;
mt19937 mersenne_engine(rd());
uniform_real_distribution<float> dist(-100.0, 100.0);
for (auto& element : matrix) {
element = dist(mersenne_engine);
}
}
int main() {
// 主机内存中的矩阵
vector<float> A(N * N), B(N * N), C(N * N);
// 使用随机数初始化矩阵
initialize_matrix(A);
initialize_matrix(B);
// 设备缓冲区用于矩阵
buffer<float, 2> bufferA(A.data(), range<2>(N, N));
buffer<float, 2> bufferB(B.data(), range<2>(N, N));
buffer<float, 2> bufferC(C.data(), range<2>(N, N));
// 使用默认选择器创建一个启用性能分析的队列
queue q{ property::queue::enable_profiling() };
// 提交一个命令组到队列以执行矩阵乘法
auto event = q.submit([&](handler& h) {
// 用于矩阵的访问器
auto accA = bufferA.get_access<access::mode::read>(h);
auto accB = bufferB.get_access<access::mode::read>(h);
auto accC = bufferC.get_access<access::mode::write>(h);
// 使用nd_range定义工作组布局
nd_range<2> work_range(range<2>(N, N), range<2>(16, 16));
// 执行矩阵乘法
h.parallel_for(work_range, [=](nd_item<2> item) {
size_t i = item.get_global_id(0);
size_t j = item.get_global_id(1);
float sum = 0.0f;
for (size_t k = 0; k < N; ++k) {
sum += accA[i][k] * accB[k][j]; // 修改索引方式
}
accC[i][j] = sum; // 修改索引方式
});
});
// 等待队列中的所有任务完成
event.wait();
// 测量内核执行所用的时间
auto start_time = event.get_profiling_info<info::event_profiling::command_start>();
auto end_time = event.get_profiling_info<info::event_profiling::command_end>();
auto execution_time = (end_time - start_time) / 1e6; // 从纳秒转换为毫秒
// 从缓冲区读回结果到主机内存
host_accessor h_accC(bufferC, read_only);
for (size_t i = 0; i < N; ++i) {
for (size_t j = 0; j < N; ++j) {
C[i * N + j] = h_accC[i][j];
}
}
// 打印结果矩阵C的第一个元素以进行验证
cout << "C[0][0] = " << C[0] << "\n";
// 打印执行时间
cout << "矩阵乘法执行时间：" << execution_time << " 毫秒\n";
return 0;
}
可以看出在上面的解决方案中，我还增加一个运行时间作为指标，因为我还想要进行比较，查看仅在CPU上进行计算和把CPU作为host，GPU作为device进行异构计算，在时间上的差异。以下是jupyterlab 上运行的结果
下面是仅仅在CPU上运行的代码：
%%writefile lab/vector_add.cpp
//==============================================================
// Copyright © Intel Corporation
//
// SPDX-License-Identifier: MIT
// =============================================================
#include <iostream>
#include <random>
#include <vector>
#include <chrono>
using namespace std;
const size_t N = 4;
// 用随机浮点数初始化矩阵的函数
void initialize_matrix(vector<float>& matrix) {
random_device rd;
mt19937 mersenne_engine(rd());
uniform_real_distribution<float> dist(-100.0, 100.0);
for (auto& element : matrix) {
element = dist(mersenne_engine);
}
}
int main() {
// 主机内存中的矩阵
vector<float> A(N * N), B(N * N), C(N * N);
// 使用随机数初始化矩阵
initialize_matrix(A);
initialize_matrix(B);
// 记录矩阵乘法开始时间
auto start = chrono::high_resolution_clock::now();
// 执行矩阵乘法
for (size_t i = 0; i < N; ++i) {
for (size_t j = 0; j < N; ++j) {
float sum = 0.0f;
for (size_t k = 0; k < N; ++k) {
sum += A[i * N + k] * B[k * N + j];
}
C[i * N + j] = sum;
}
}
// 记录矩阵乘法结束时间
auto end = chrono::high_resolution_clock::now();
// 计算并打印执行时间
auto duration = chrono::duration_cast<chrono::microseconds>(end - start).count();
cout << "矩阵乘法执行时间：" << duration << " 微秒\n";
// 打印结果矩阵C的第一个元素以进行验证
cout << "C[0][0] = " << C[0] << "\n";
return 0;
}
我们可以看到，相比仅在CPU上运行，使用异构计算明显加快了运算速度。
下面我们来看任务三：
（4）任务3：统一共享内存（简称“USM”）
统一共享内存在主机和设备（GPU）之间提供了一个统一的存储模式。修改你的程序以使用统一共享内存来实现内存分配和数据转换，从而替代缓存和存储。如果使用不同种类的USM将获得加分。重写代码以使用统一共享内存来实现内存分配和数据转换（如需要），照常运行Kernel来验证程序运行结果。
在任务三中，我们将使用不同的USM（host、shared、device），并用运行时间作为指标，编写程序，发现他们的差异，并进行原因分析。
a.host_USM
难点分析
- 内存分配与释放：使用
malloc_host在主机上分配内存，并确保在程序结束时正确释放这些内存。 - 内存访问与同步：虽然使用主机USM可以简化内存访问，但需要确保内核执行结束后再从主机端访问数据。
关键代码：
1.内存分配
float* A = malloc_host<float>(N * N, default_selector{});
float* B = malloc_host<float>(N * N, default_selector{});
float* C = malloc_host<float>(N * N, default_selector{});
2.矩阵乘法执行
q.submit([&](handler& h) {
h.parallel_for(range<2>(N, N), [=](id<2> id) {
size_t i = id[0];
size_t j = id[1];
float sum = 0.0f;
for (size_t k = 0; k < N; ++k) {
sum += A[i * N + k] * B[k * N + j];
}
C[i * N + j] = sum;
});
}).wait();
3.不要忘了内存释放
free(A, q.get_context());
free(B, q.get_context());
free(C, q.get_context());
完整代码：
#include <CL/sycl.hpp>
#include <iostream>
#include <random>
#include <chrono>
using namespace sycl;
using namespace std;
const size_t N = 1024;
void initialize_matrix(float* matrix, size_t size) {
random_device rd;
mt19937 mersenne_engine(rd());
uniform_real_distribution<float> dist(-100.0, 100.0);
for (size_t i = 0; i < size; ++i) {
matrix[i] = dist(mersenne_engine);
}
}
int main() {
queue q{};
// 使用 queue 的 context 来分配内存
float* A = malloc_host<float>(N * N, q.get_context());
float* B = malloc_host<float>(N * N, q.get_context());
float* C = malloc_host<float>(N * N, q.get_context());
initialize_matrix(A, N * N);
initialize_matrix(B, N * N);
auto start = std::chrono::high_resolution_clock::now();
q.submit([&](handler& h) {
h.parallel_for(range<2>(N, N), [=](id<2> id) {
size_t i = id[0];
size_t j = id[1];
float sum = 0.0f;
for (size_t k = 0; k < N; ++k) {
sum += A[i * N + k] * B[k * N + j];
}
C[i * N + j] = sum;
});
}).wait();
auto end = std::chrono::high_resolution_clock::now();
auto duration = std::chrono::duration_cast<std::chrono::milliseconds>(end - start).count();
cout << "C[0][0] = " << C[0] << "\n";
cout << "使用主机分配的USM进行矩阵乘法的执行时间：" << duration << " 毫秒\n";
free(A, q.get_context());
free(B, q.get_context());
free(C, q.get_context());
return 0;
}
运行结果：
b.改成使用shared_USM
#include <CL/sycl.hpp>
#include <iostream>
#include <random>
#include <chrono>
using namespace sycl;
using namespace std;
const size_t N = 1024;
void initialize_matrix(float* matrix, size_t size) {
random_device rd;
mt19937 mersenne_engine(rd());
uniform_real_distribution<float> dist(-100.0, 100.0);
for (size_t i = 0; i < size; ++i) {
matrix[i] = dist(mersenne_engine);
}
}
int main() {
// 创建队列，使用default_selector_v和性能分析属性
queue q(default_selector_v, property::queue::enable_profiling());
float* A = malloc_shared<float>(N * N, q);
float* B = malloc_shared<float>(N * N, q);
float* C = malloc_shared<float>(N * N, q);
initialize_matrix(A, N * N);
initialize_matrix(B, N * N);
auto start = std::chrono::high_resolution_clock::now();
q.submit([&](handler& h) {
h.parallel_for(range<2>(N, N), [=](id<2> id) {
size_t i = id[0];
size_t j = id[1];
float sum = 0.0f;
for (size_t k = 0; k < N; ++k) {
sum += A[i * N + k] * B[k * N + j];
}
C[i * N + j] = sum;
});
}).wait();
auto end = std::chrono::high_resolution_clock::now();
auto duration = std::chrono::duration_cast<std::chrono::milliseconds>(end - start).count();
cout << "C[0][0] = " << C[0] << "\n";
cout << "使用Shared-USM进行矩阵乘法的执行时间：" << duration << " 毫秒\n";
free(A, q);
free(B, q);
free(C, q);
return 0;
}
运行结果：
c.device_USM
#include <CL/sycl.hpp>
#include <iostream>
#include <random>
#include <chrono>
using namespace sycl;
using namespace std;
const size_t N = 1024;
void initialize_matrix(float* matrix, size_t size, queue& q) {
float* temp_matrix = malloc_host<float>(size, q.get_context());
random_device rd;
mt19937 mersenne_engine(rd());
uniform_real_distribution<float> dist(-100.0, 100.0);
for (size_t i = 0; i < size; ++i) {
temp_matrix[i] = dist(mersenne_engine);
}
q.memcpy(matrix, temp_matrix, sizeof(float) * size).wait();
free(temp_matrix, q.get_context());
}
int main() {
queue q(default_selector{}, property::queue::enable_profiling());
float* A = malloc_device<float>(N * N, q);
float* B = malloc_device<float>(N * N, q);
float* C = malloc_device<float>(N * N, q);
initialize_matrix(A, N * N, q);
initialize_matrix(B, N * N, q);
auto start = std::chrono::high_resolution_clock::now();
q.submit([&](handler& h) {
h.parallel_for(range<2>(N, N), [=](id<2> id) {
size_t i = id[0];
size_t j = id[1];
float sum = 0.0f;
for (size_t k = 0; k < N; ++k) {
sum += A[i * N + k] * B[k * N + j];
}
C[i * N + j] = sum;
});
}).wait();
auto end = std::chrono::high_resolution_clock::now();
auto duration = std::chrono::duration_cast<std::chrono::milliseconds>(end - start).count();
float* C_host = malloc_host<float>(N * N, q.get_context());
q.memcpy(C_host, C, sizeof(float) * N * N).wait();
cout << "C[0][0] = " << C_host[0] << "\n";
cout << "使用Device-Allocated USM进行矩阵乘法的执行时间：" << duration << " 毫秒\n";
free(A, q);
free(B, q);
free(C, q);
free(C_host, q.get_context());
return 0;
}
运行结果
差异因素分析：
1. Device-Allocated USM : 时间最短。
原因分析：
高带宽和并行性：GPU通常具有非常高的内存带宽，可以同时处理大量的操作，特别是对于像矩阵乘法这样的并行计算密集型任务。
近距离内存访问：在设备上分配的内存可以被GPU直接访问，没有跨设备的数据传输开销。
优化的内存访问：使用设备内存可能意味着更优化的内存访问模式，减少了内存访问延迟。
2. Host-Allocated USM : 时间最长。
原因分析：
数据传输开销：虽然主机分配的内存可以被GPU访问，但可能涉及显式的数据传输，会增加额外的延迟。
可能的同步问题：在某些情况下，主机和设备之间可能需要同步，这会导致额外的性能开销。
内存访问速度：主机内存的带宽通常低于GPU内存，且访问速度可能慢于设备上直接分配的内存。
3. Shared-Allocated USM: 时间介于两者之间。
原因分析：
平衡的内存访问：共享内存旨在平衡主机和设备的内存访问需求，提供了一种折中的性能。
潜在的优化：可能在某些硬件上进行了优化，以提高对共享内存的访问速度。
减少数据迁移：减少了与主机分配的USM相比的数据迁移需求，但仍可能面临一些内存访问开销。
综合来看，影响运算性能（这里选取运算时间作为指标）的因素有：
内存带宽：内存与计算单元之间数据传输的速度。
内存访问模式：缓存的使用效率，以及访问模式是否能够充分利用硬件的内存访问优化。
数据传输：数据在主机和设备之间传输的需要，包括数据传输的次数和量。
硬件优化：特定硬件架构的优化，如GPU的内存访问优化，可以大大影响性能。
同步机制：主机和设备之间同步数据的机制，可能会增加延迟。
并行计算资源：GPU拥有的并行计算资源远多于CPU，对于并行计算任务能提供更好的性能。
四、实验中的挑战与心得体会
挑战一：理解SYCL编程模型
作为oneAPI的关键部分，SYCL为我提供了一个充满挑战的学习曲线。刚开始接触SYCL时，理解其异构编程概念和抽象级别对我来说非常困难。例如，缓冲区、访问器、队列和内核的概念及其在并行计算中的协同作用，初看起来令人望而却步。为了克服这一挑战，我深入研究了SYCL官方文档和oneAPI相关教程。通过反复练习编写简单程序并逐渐增加复杂度，我逐渐掌握了SYCL的核心原理。此外，参与在线社区的讨论也大大帮助了我解决了理解过程中的疑问。
挑战二：性能优化
在oneAPI的性能优化阶段，我意识到，单纯将计算任务迁移到GPU并不意味着能自动获得最佳性能。SYCL提供的nd_range配置和内存访问模式调整为性能优化提供了更多可能性。我通过实验不同的工作组配置，优化了矩阵乘法内核的性能。通过分析内核的内存访问模式，我学会了如何调整内存访问策略，以减少内存延迟和提高带宽利用率。
挑战三：使用不同的USM类型
在oneAPI的框架下，我被要求使用不同的统一共享内存（USM）类型执行矩阵乘法，并比较它们的性能。每种USM类型（如显式USM、隐式USM）都有其独特的优势和局限性。理解这些类型的高效使用方式是一个挑战。我通过编写多个版本的矩阵乘法程序，每个版本采用不同的USM策略，深入探究了它们的性能差异。通过分析和比较，我不仅理解了不同USM类型对性能的影响，还学会了根据特定的计算任务选择最合适的USM类型。
心得体会
通过这个实验，我不仅学会了使用SYCL进行高效的异构编程，还深刻理解了性能优化在实际应用中的重要性。我发现，即使是看似简单的矩阵乘法，也存在许多优化的空间。这个实验过程让我意识到，深入理解硬件架构和编程模型对于进行有效的并行编程至关重要。
最重要的是，这个实验提高了我的问题分析和解决能力。面对挑战时，我学会了不断尝试不同解决方案，直至找到最佳方法。我不仅提升了自己的编程技能，也增强了自我学习和适应新技术的能力。
扫一扫
558
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
