0
1
0
0
K100 DCU的OpenMP与OpenACC性能问题
2025-03-18 10:54:39
135
#OpenMP# #OpenACC#
用Element-wise乘法为例，基于K100 DCU写了OpenMP与OpenACC版本，加速效果在2x至4x左右，加速比过低。
2、内容
操作系统版本：CentOS Linux release 7.6.181
DTK版本：DTK24.04.1
驱动版本：/
DCU类型：K100
问题描述：编译执行OpenMP与OpenACC版本，性能加速只有2x至4x，加速比过低。请问代码或编译是否有问题？还是访存带宽原因？
运行结果：
CPU time: 0.295973s, GFLOPS: 0.22674
OpenACC time: 0.523789s, GFLOPS: 0.128122
Verification: PASSED
CPU time: 0.112369s, GFLOPS: 0.597219
OpenMP time: 0.381527s, GFLOPS: 0.175896
Verification: PASSED
--
楼主
0 点赞
0 帖子
楼主热帖
暂无数据常见问题新帖
暂无数据
矩阵逐元素相乘好像是这样的，GFLOPS不高，性能瓶颈都在带宽上，加速比也不会很好。建议可以试试矩阵乘法，比如GNU的这个示例代码：https://gcc.gnu.org/wiki/OpenACC/OpenACC%20%27kernels%27%20Construct%20Optimization%20Tutorial
看看加速比会不会好一点#include <iostream>
#include <vector>
#include <chrono>
#include <cstdlib>
#ifdef OPENMP_ENABLED
#include <omp.h>
#endif
#ifdef OPENACC_ENABLED
#include <openacc.h>
#endif
// Element-wise 乘法：C = A ⊙ B
void elementwise_cpu(const std::vector<double> &A,
const std::vector<double> &B,
std::vector<double> &C,
int m, int n)
{
#pragma omp parallel for collapse(2) if(m*n > 10000)
for (int i = 0; i < m; ++i) {
for (int j = 0; j < n; ++j) {
C[i*n + j] = A[i*n + j] * B[i*n + j];
}
}
}
#ifdef OPENMP_ENABLED
// void elementwise_omp(const std::vector<double> &A,
// const std::vector<double> &B,
// std::vector<double> &C,
// int m, int n)
// {
// auto A_ptr = A.data();
// auto B_ptr = B.data();
// auto C_ptr = C.data();
// #pragma omp target data map(to: A_ptr[0:m*n], B_ptr[0:m*n]) map(from: C_ptr[0:m*n])
// {
// #pragma omp target teams distribute parallel for collapse(2)
// for(int i = 0; i < m; ++i) {
// for(int j = 0; j < n; ++j) {
// C_ptr[i*n + j] = A_ptr[i*n + j] * B_ptr[i*n + j];
// }
// }
// }
// }
void elementwise_omp(const std::vector<double> &A,
const std::vector<double> &B,
std::vector<double> &C,
int m, int n)
{
auto A_ptr = A.data();
auto B_ptr = B.data();
auto C_ptr = C.data();
#pragma omp target teams distribute parallel for collapse(2) map(to: A_ptr[0:m*n], B_ptr[0:m*n]) map(from: C_ptr[0:m*n])
for(int i = 0; i < m; ++i) {
for(int j = 0; j < n; ++j) {
C_ptr[i*n + j] = A_ptr[i*n + j] * B_ptr[i*n + j];
}
}
}
#endif
#ifdef OPENACC_ENABLED
void elementwise_acc(const std::vector<double> &A,
const std::vector<double> &B,
std::vector<double> &C,
int m, int n)
{
auto A_ptr = A.data();
auto B_ptr = B.data();
auto C_ptr = C.data();
#pragma acc data copyin(A_ptr[0:m*n], B_ptr[0:m*n]) copyout(C_ptr[0:m*n])
{
// #pragma acc parallel loop gang vector collapse(2)
const int VECTOR_LENGTH = 256; // 海光DCU推荐值
#pragma acc parallel loop gang vector vector_length(VECTOR_LENGTH) independent
for (int i = 0; i < m; ++i) {
for (int j = 0; j < n; ++j) {
C_ptr[i*n + j] = A_ptr[i*n + j] * B_ptr[i*n + j];
}
}
}
}
#endif
// 结果校验
bool check_result(const std::vector<double> &C1, const std::vector<double> &C2)
{
if (C1.size() != C2.size())
{
std::cerr << "Error: Result vectors have different sizes." << std::endl;
return false;
}
for (size_t i = 0; i < C1.size(); ++i)
{
if (std::abs(C1[i] - C2[i]) > 1e-5)
{
std::cerr << "Error: Mismatch at index " << i << ". Expected: " << C1[i] << ", Actual: " << C2[i] << std::endl;
return false;
}
}
return true;
}
// 计算强度调整为 element-wise 操作
double calculate_flops(int m, int n) {
return 1.0 * m * n; // 每个元素1次乘法
}
int main(int argc, char *argv[])
{
if (argc != 3) { // 参数减少为 m 和 n
std::cerr << "Usage: " << argv[0] << " m n" << std::endl;
return 1;
}
int m = std::atoi(argv[1]);
int n = std::atoi(argv[2]);
// 矩阵维度统一为 m x n
std::vector<double> A(m * n);
std::vector<double> B(m * n);
std::vector<double> C_cpu(m * n);
std::vector<double> C_parallel(m * n);
// 初始化矩阵（随机值0~1）
#pragma omp parallel for
for (int i = 0; i < m*n; ++i) {
A[i] = static_cast<double>(rand()) / RAND_MAX;
B[i] = static_cast<double>(rand()) / RAND_MAX;
}
// CPU 版本基准测试
auto start_cpu = std::chrono::high_resolution_clock::now();
elementwise_cpu(A, B, C_cpu, m, n);
auto end_cpu = std::chrono::high_resolution_clock::now();
std::chrono::duration<double> duration_cpu = end_cpu - start_cpu;
double gflops_cpu = calculate_flops(m, n) / duration_cpu.count() / 1e9;
std::cout << "CPU time: " << duration_cpu.count()
<< "s, GFLOPS: " << gflops_cpu << std::endl;
// 加速版本（OpenMP/OpenACC）
auto start_parallel = std::chrono::high_resolution_clock::now();
#ifdef OPENMP_ENABLED
elementwise_omp(A, B, C_parallel, m, n);
#elif defined(OPENACC_ENABLED)
elementwise_acc(A, B, C_parallel, m, n);
#endif
auto end_parallel = std::chrono::high_resolution_clock::now();
std::chrono::duration<double> duration_parallel = end_parallel - start_parallel;
double gflops_parallel = calculate_flops(m, n) / duration_parallel.count() / 1e9;
#ifdef OPENMP_ENABLED
std::cout << "OpenMP time: " << duration_parallel.count()
<< "s, GFLOPS: " << gflops_parallel << std::endl;
#elif defined(OPENACC_ENABLED)
std::cout << "OpenACC time: " << duration_parallel.count()
<< "s,