DPC++实验---矩阵乘法示例

最新推荐文章于 2024-10-31 18:26:46 发布
最新推荐文章于 2024-10-31 18:26:46 发布
阅读量178 收藏 2
点赞数
文章标签： c++ 矩阵 数学建模
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/weixin_62678986/article/details/131269472
版权
文章通过两个练习展示了使用SYCL进行GPU和CPU上的矩阵乘法运算。代码实现了GPU和CPU上的kernel，对比了不同参数设置（如tile尺寸）对运算性能的影响。实验发现，改变tile尺寸对GPU执行时间有显著影响，而对CPU执行时间影响较小。同时，当tile尺寸过大时，可能导致执行时间增加甚至出错。
摘要由CSDN通过智能技术生成

目标：

        实现代码实践运行并分析结果

内容：

一、练习一

代码如下：

#include <chrono>
#include <iostream>
#include <CL/sycl.hpp>
#define random_float() (rand() / double(RAND_MAX))
using namespace std;
using namespace sycl;
// return execution time
double gpu_kernel(float *A, float *B, float *C, int M, int N, int K, int block_size, sycl::queue &q) {
  // define the workgroup size and mapping
  auto grid_rows = (M + block_size - 1) / block_size * block_size;
  auto grid_cols = (N + block_size - 1) / block_size * block_size;
  auto local_ndrange  = range<2>(block_size, block_size);
  auto global_ndrange = range<2>(grid_rows, grid_cols);
  double duration = 0.0f;
  auto e = q.submit([&](sycl::handler &h) {
      h.parallel_for<class k_name_t>(
          sycl::nd_range<2>(global_ndrange, local_ndrange), [=](sycl::nd_item<2> index) {
              int row = index.get_global_id(0);
              int col = index.get_global_id(1);
              float sum = 0.0f;
              for (int i = 0; i < K; i++) {
                sum += A[row * K + i] * B[i * N  + col];
              }
              C[row * N + col] = sum;
          });
    });
    e.wait();
    duration += (e.get_profiling_info<info::event_profiling::command_end>() -
    e.get_profiling_info<info::event_profiling::command_start>()) /1000.0f/1000.0f;
    return(duration);
}
// return execution time
double cpu_kernel(float *cA, float *cB, float *cC, int M, int N, int K) {
    double duration = 0.0;
    std::chrono::high_resolution_clock::time_point s, e;
    // Single Thread Computation in CPU 
    s = std::chrono::high_resolution_clock::now();
    for(int i = 0; i < M; i++) {
        for(int j = 0; j < N; j++) {
            float sum = 0.0f;
            for(int k = 0; k < K; k++) {
                sum +=  cA[i * K + k] * cB[k * N  + j];
            }
            cC[i * N + j] = sum;
        }
    }
    e = std::chrono::high_resolution_clock::now();
    duration = std::chrono::duration<float, std::milli>(e - s).count();
    return(duration);
}
int verify(float *cpu_res, float *gpu_res, int length){
    int err = 0;
    for(int i = 0; i < length; i++) {
       if( fabs(cpu_res[i] - gpu_res[i]) > 1e-3) {
          err++;
          printf("\n%lf, %lf", cpu_res[i], gpu_res[i]);
       } 
    }
    return(err);
}
int gemm(const int M, 
         const int N, 
         const int K, 
         const int block_size,
         const int iterations, 
         sycl::queue &q) {
  cout << "Problem size: c(" << M << "," <<  N << ") ="
       << " a(" << M << "," << K << ") *" 
       << " b(" << K << "," << N << ")\n";
 
  auto A = malloc_shared<float>(M * K, q);
  auto B = malloc_shared<float>(K * N, q);
  auto C = malloc_shared<float>(M * N, q);
  auto C_host = malloc_host<float>(M * N, q);
  // init the A/B/C
  for(int i=0; i < M * K; i++) {
      A[i] = random_float();
  }
  for(int i=0; i < K * N; i++) {
      B[i] = random_float();
  }
  for(int i=0; i < M * N; i++) {
      C[i] = 0.0f;
      C_host[i] = 0.0f;
  }
  double flopsPerMatrixMul
      = 2.0 * static_cast<double>(M) * static_cast<double>(N) * static_cast<double>(K);
  double duration_gpu = 0.0f;
  double duration_cpu = 0.0f;
  // GPU compuation and timer 
  int warmup = 10;
  for (int run = 0; run < iterations + warmup; run++) {
    float duration = gpu_kernel(A, B, C, M, N, K, block_size, q);
    if(run >= warmup) duration_gpu += duration;
  }
  duration_gpu = duration_gpu / iterations;
  // CPU compuation and timer 
  warmup = 2;
  for(int run = 0; run < iterations/2 + warmup; run++) {
      float duration = cpu_kernel(A, B, C_host, M, N, K);
      if(run >= warmup) duration_cpu += duration;
  }
  duration_cpu = duration_cpu / iterations/2;
  // Compare the resutls of CPU and GPU 
  int errCode = 0;
  errCode = verify(C_host, C, M*N);
  if(errCode > 0) printf("\nThere are %d errors\n", errCode);
  printf("\nPerformance Flops = %lf, \n" 
          "GPU Computation Time = %lf (ms); \n"
          "CPU Computaiton Time = %lf (ms); \n", 
          flopsPerMatrixMul, duration_gpu, duration_cpu);
 
  free(A, q);
  free(B, q);
  free(C, q);
  free(C_host, q);
  return(errCode);
}
 
int main() {
  auto propList = cl::sycl::property_list {cl::sycl::property::queue::enable_profiling()};
  queue my_gpu_queue( cl::sycl::gpu_selector_v , propList);
  int errCode = gemm(2000, 2000, 2000, 4, 10, my_gpu_queue);
  return(errCode);
}

可通过修改MNK值决定矩阵参数，实现不同结果进行比较，此例设置M=N=K=2000

运行结果：

二、练习二

代码如下：

#include <chrono>
#include <iostream>
#include <CL/sycl.hpp>
#define random_float() (rand() / double(RAND_MAX))
using namespace std;
using namespace sycl;
#define tileY 2 
#define tileX 2
// return execution time
double gpu_kernel(float *A, float *B, float *C, 
                  int M, int N, int K, 
                  int BLOCK, sycl::queue &q) {
  // define the workgroup size and mapping
  auto grid_rows = M / tileY;
  auto grid_cols = N / tileX;
  auto local_ndrange  = range<2>(BLOCK, BLOCK);
  auto global_ndrange = range<2>(grid_rows, grid_cols);
  double duration = 0.0f;
  auto e = q.submit([&](sycl::handler &h) {
      h.parallel_for<class k_name_t>(
          sycl::nd_range<2>(global_ndrange, local_ndrange), [=](sycl::nd_item<2> index) {
              int row = tileY * index.get_global_id(0);
              int col = tileX * index.get_global_id(1);
              float sum[tileY][tileX] = {0.0f};
              float subA[tileY] = {0.0f};
              float subB[tileX] = {0.0f};
               // core computation
              for (int k = 0; k < N; k++) {
                // read data to register
                for(int m = 0; m < tileY; m++) {
                    subA[m] = A[(row + m) * N + k];
                } 
                for(int p = 0; p < tileX; p++) {
                    subB[p] = B[k * N + p + col];
                } 
                for (int m = 0; m < tileY; m++) {
                  for (int p = 0; p < tileX; p++) {
                    sum[m][p] += subA[m] * subB[p];
                  }
                }
              } //end of K
              // write results back
              for (int m = 0; m < tileY; m++) {
                for (int p = 0; p < tileX; p++) {
                  C[(row + m) * N + col + p] = sum[m][p];
                }
              }
          });
    });
    e.wait();
    duration += (e.get_profiling_info<info::event_profiling::command_end>() -
    e.get_profiling_info<info::event_profiling::command_start>()) /1000.0f/1000.0f;
    return(duration);
}
// return execution time
double cpu_kernel(float *cA, float *cB, float *cC, int M, int N, int K) {
    double duration = 0.0;
    std::chrono::high_resolution_clock::time_point s, e;
    // Single Thread Computation in CPU 
    s = std::chrono::high_resolution_clock::now();
    for(int i = 0; i < M; i++) {
        for(int j = 0; j < N; j++) {
            float sum = 0.0f;
            for(int k = 0; k < K; k++) {
                sum +=  cA[i * K + k] * cB[k * N  + j];
            }
            cC[i * N + j] = sum;
        }
    }
    e = std::chrono::high_resolution_clock::now();
    duration = std::chrono::duration<float, std::milli>(e - s).count();
 
    return(duration);
}
int verify(float *cpu_res, float *gpu_res, int length){
    int err = 0;
    for(int i = 0; i < length; i++) {
       if( fabs(cpu_res[i] - gpu_res[i]) > 1e-3) {
          err++;
          printf("\n%lf, %lf", cpu_res[i], gpu_res[i]);
       } 
    }
    return(err);
}
int gemm(const int M, 
         const int N, 
         const int K, 
         const int block_size,
         const int iterations, 
         sycl::queue &q) {
  cout << "Problem size: c(" << M << "," <<  N << ") ="
       << " a(" << M << "," << K << ") *" 
       << " b(" << K << "," << N << ")\n";
  auto A = malloc_shared<float>(M * K, q);
  auto B = malloc_shared<float>(K * N, q);
  auto C = malloc_shared<float>(M * N, q);
  auto C_host = malloc_host<float>(M * N, q);
  // init the A/B/C
  for(int i=0; i < M * K; i++) {
      A[i] = random_float();
  }
  for(int i=0; i < K * N; i++) {
      B[i] = random_float();
  }
  for(int i=0; i < M * N; i++) {
      C[i] = 0.0f;
      C_host[i] = 0.0f;
  }
  double flopsPerMatrixMul
      = 2.0 * static_cast<double>(M) * static_cast<double>(N) * static_cast<double>(K);
  double duration_gpu = 0.0f;
  double duration_cpu = 0.0f;
  // GPU compuation and timer 
  int warmup = 10;
  for (int run = 0; run < iterations + warmup; run++) {
    float duration = gpu_kernel(A, B, C, M, N, K, block_size, q);
    if(run >= warmup) duration_gpu += duration;
  }
  duration_gpu = duration_gpu / iterations;
  // CPU compuation and timer 
  warmup = 2;
  for(int run = 0; run < iterations/2 + warmup; run++) {
      float duration = cpu_kernel(A, B, C_host, M, N, K);
      if(run >= warmup) duration_cpu += duration;
  }
  duration_cpu = duration_cpu / iterations/2;
  // Compare the resutls of CPU and GPU 
  int errCode = 0;
  errCode = verify(C_host, C, M*N);
  if(errCode > 0) printf("\nThere are %d errors\n", errCode);
  printf("\nGEMM size M = %d, N = %d, K = %d", M, N, K);
  printf("\nWork-Group size = %d * %d, tile_X = %d, tile_Y = %d", block_size, block_size, tileX, tileY);
  printf("\nPerformance Flops = %lf, \n" 
          "GPU Computation Time = %lf (ms); \n"
          "CPU Computaiton Time = %lf (ms); \n", 
          flopsPerMatrixMul, duration_gpu, duration_cpu);
 
  free(A, q);
  free(B, q);
  free(C, q);
  free(C_host, q);
  return(errCode);
}
int main() {
  auto propList = cl::sycl::property_list {cl::sycl::property::queue::enable_profiling()};
  queue my_gpu_queue( cl::sycl::gpu_selector_v , propList);
  int errCode = gemm(512, 512, 512, /* GEMM size, M, N, K */
                     4,             /* workgroup size */ 
                     10,            /* repeat time */   
                     my_gpu_queue);
  return(errCode);
}

测试不同tile_X和tile_Y大小下（ tile_X和tile_Y可以不一致），矩阵计算的性能

运行结果：

（2,2）

（2,4）

（2,8）

（4,2）

（8,2）

 分析：

1.在tileX，tileY的变化中，可以看到CPU执行时间基本没有改变，说明其变化对于CPU来说并无太大影响。

2.保持tileX不变，当tileY增大时，GPU执行时间有所改变，但改变较小，说明其改变对GPU执行时间影响较小；保持tileY不变，当tileX增大时，GPU执行时间变化较大，说明tileX的变化对GPU执行时间影响很大。

3.当tileX和tileY达到8以上时GPU和CPU执行时间有很大概率爆掉，出现错误现象，下面列出两组错误数据：

 

slagtall
关注
通过DPC ++从CUDA到SYCL 2020进行光线追踪
04-11
对于示例中的代码，`simple_raytracing_cuda-main.zip`包含了原始的CUDA光线追踪实现，而`simple_raytracing_sycl2020-main.zip`是转换后的DPC++/SYCL版本。阅读和比较这两个版本可以帮助理解转换过程。`Raytracing-...
SHU--算法设计实验二：矩阵连乘问题
dxxmsl的博客
11-11 1099
在这次实验的过程中，如何得知是哪个地方要加括号这个问题我思考了很久，在一遍遍理清楚整个算法的过程后，我有了初步的了解，加括号的位置其实就是k 的对应序号的矩阵，在写算法时我们就可以用另外的数组记录下对应位置的k值。在分析问题的最优子结构性质时，所用的方法具有普遍性：首先假设由问题的最优解导出的子问题的解不是最优的，然后再设法说明在这个假设下可构造出比原问题最优解更好的解，从而导致矛盾。设计算A[i:j]（矩阵A从i乘到j），1≤i≤j≤n，所需要的最少数乘次数m[i,j]，则原问题的最优值为m[1,n]。
超微 X12DPC-6主板用户手册
03-25
### 超微 X12DPC-6 主板用户手册关键知识点解析 #### 一、产品概述 超微 X12DPC-6 主板是超微计算机股份有限公司（以下简称“Supermicro”）推出的一款高性能服务器主板。该主板的设计旨在为用户提供稳定可靠的...
DPC.zip_DPC_RP-DPC MATLAB代码_dpc算法_dpc算法 matlab_数据集
07-14
matalab个人实现的DPC算法，外加数据集
oneAPI-samples:英特尔oneAPI工具包示例
03-19
英特尔oneAPI工具包示例 oneAPI样本存储库提供了Intel oneAPI工具箱的代码样本。 我们建议签出存储库的特定发行版本。。 不能保证master分支上的代码示例的最新版本是稳定的。 代码样例 ...../DPC+
DPC-SLIC实验报告1
08-08
【DPC-SLIC实验报告1】 在计算机视觉和图像处理领域，聚类是一种常见的无监督学习方法，用于发现数据中的内在结构和模式。本实验报告主要探讨了使用Superpixel算法SLIC（Simple Linear Iterative Clustering）以及...
《C++显式类型转换：解析多种转换方式的奥秘》
xy520521的博客
10-27 699
当我们从一个存储 Shape 指针的容器中取出指针并希望将其转换为具体的派生类指针以调用特定的绘制方法时，动态_cast可以确保转换的安全性，避免在对象类型不匹配时出现错误。需要注意的是，使用 const_cast修改一个原本被定义为 const 的对象可能会导致未定义的行为，因为这违反了 const 的语义。例如，当我们有一个 const 引用指向一个对象，但在某个特定的函数内部，我们知道这个函数不会修改对象的值，但需要将其传递给一个非 const 版本的函数时，可以使用 const_cast。
「C/C++」C/C++标准库之#include＜cstdlib＞通用工具库
何曾参静谧的博客
10-28 576
是C++标准库中的一个头文件，它提供了各种通用工具函数，这些函数最初来源于C标准库中的，并在C++中进行了标准化和扩展。在C++编程中，头文件具有重要地位，它涵盖了字符串转换、类型转换、随机数生成、内存管理和系统调用等多个方面的功能。通过这些实例，我们可以看到头文件在C++编程中的广泛应用。它提供了一系列功能强大的函数，使得开发者能够方便地处理字符串转换、随机数生成、内存管理等任务。在实际开发中，合理利用这些函数可以大大提高编程效率和代码质量。
C/C++：WinSock_TCP通信实例
qq_55038440的博客
10-28 977
1.AF地址族 1.SOCKADDR本地地址结构体 2.SOCKADDR_IN TCP/IP地址结构体 案例：指定端口号代码 案例：指定 IP地址代码 三、监听套接字 四、接收连接请求 1.accept()函数 函数原型如下： 2.connect()函数 函数原型如下： 五、收发数据 1.send()函数 函数原型如下 2.recv()函数 函数原型如下
第九天: 深入探索C语言中的动态内存管理
最新发布
qq_57484399的博客
10-31 162
简单来说，就是C语言没有直接提供给我们的数据类型，但是C语言给我们留了一个后门，允许我们用一些特殊的关键字来定义自己的数据类型。枚举类型允许我们定义一组相关的常量，这些常量就是枚举类型的取值。听起来有点绕？我给你举个例子。比如说，我们要表示一周的七天，我们可以用字符串来表示，但处理起来比较麻烦。用整数表示就简单多了，星期一是1，星期二是2，以此类推。这样的变量就可以定义为枚举类型，它的取值范围就是1到7。首先，我们用struct关键字来定义一个结构体。
C++ 优先算法——复写零（双指针）
2301_80873544的博客
10-28 631
这个也要用双指针算法来去寻找它，我们把dest定义为最后需要复写的位置（结果中数组最后一个位置），因此刚开始时，并不知道最后位置在哪，让dest指向-1；cur的作用：从前往后遍历数组，决定dest走1步还是2步。所以可以仅需要判断dest是否走向最后位置来判断是否找到最后要复写的数。// 先找到最后一个要复写的数// 复写位置// 遍历数组if (arr[cur] == 0) { //遇到0，dest移动两位dest += 2;
C++二级 字符串判断是否是子列
xzal12的博客
10-28 212
C++二级 字符串判断是否是子列
C++--this关键字
qq_63152518的博客
10-26 1372
this 的类型是 ClassName* const，其中 ClassName 是类的名称。const 表示指针本身是常量，意味着你不能改变 this 指针指向的地址（即它始终指向调用该函数的对象）。
【C++】动态库动态加载实例详解
richard_yuu的博客
10-26 979
DLL动态库是一种包含程序代码和数据的可执行文件，它可以在运行时被程序加载和链接。通过将功能封装成DLL，我们可以实现代码的模块化封装，使程序更加灵活和易于维护。此外，为了方便第三方厂商接入我们的系统进行二次开发，我们通常会将相关业务封装成SDK动态库供他们使用。在C++程序中，通常会使用多个DLL库。在发布程序时，这些DLL库需要被打包到安装包中一起发布。安装过程中，这些DLL会被放置到程序的安装目录中（有时个别库可能需要放到系统路径中）。
STL动态数组
dazys的博客
10-26 492
例如，对于容器 vector v，可以使用 v[index]来访问它的第 index 个元素。其中，0≤index≤v.size()-1，v.size()表示 vector 中元素的个数。迭代器类似于指针，指向vector中元素的位置，可以使用迭代器来访问vector中的元素。a.end() // 指向数组a的结束位置的下一个位置。a.begin() // 指向数组a的开始位置。1.容器：容纳包含一组元素的对象。2.迭代器：提供访问容器的方法。// t1指向数组a的开始位置。// t2指向数组a结束位置。
【每日C/C++问题】
2301_80313293的博客
10-26 986
do while 与 while 基本相同，表达式为真才会执行循环语句。顺序结构：从第一条语句到最后一条语句完全按顺序执行，是简单的顺序结构；选择结构：通过某个给定条件进行判断，条件为真或假时分别执行不同的程序内容,，后面语句与if表达式判断无关。如果想当if表达式为真时执行多条语句，选择结构：按照给定条件有选择地执行程序中的语句。表达式1：初始化部分，用于初始化循环变量，当表达式为真时执行循环语句，然后再次判断…表达式2：条件判断部分，用于判断循环终止。表达式3：调整部分，用于循环条件的调整。
C++二级2023.9题及答案 -- 部分题
xzal12的博客
10-27 325
C++二级2023.9题及答案 -- 部分题
C++ 模板
2301_80758704的博客
10-28 667
class 类模板名// 类内成员定义// 动态顺序表// 注意：Vector不是具体的类，是编译器根据被实例化的类型生成具体类的模具public :, _size(0){}// 使用析构函数演示：在类中声明，在类外定义。~Vector();// 注意：类模板中函数放在类外进行定义时，需要加模板参数列表if(_pData)
Python与C++的区别：你适合哪种编程语言？
邓校长的编程课堂专栏
10-27 634
或希望在数据科学、人工智能、Web开发等领域快速应用编程，
DPC++ linux 环境变量
09-22
DPC在Linux环境中，需要设置DM_HOME环境变量和动态链接库文件的加载路径。可以通过以下步骤进行配置： 1. 切换至dmdba用户：`su - dmdba` 2. 执行命令使环境变量生效：`source .bash_profile` 3. 切换到root用户并...
评论 1
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值