DLUT操作系统上机补充实验5——DPC++矩阵乘法

最新推荐文章于 2023-06-17 23:56:37 发布
最新推荐文章于 2023-06-17 23:56:37 发布
阅读量514 收藏 4
点赞数 3
文章标签： python 开发语言
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/qq_64080447/article/details/131095041
版权

书接上回，上一次实验完成了几个实验用例在本地和云端的编译，但云端始终无法调用gpu_selector_v（可能是我自己的问题），后续尝试在本地可以使用gpu_selector_v。

这次的矩阵乘法实验选需要比较cpu时间与gpu时间，故选择在本地完成实验。

实验课练习1：

代码：

（来源：oneAPI_course/gemm_basic.cpp at main · pengzhao-intel/oneAPI_course · GitHub）（经过略微修改，将gpu_selector_{} 改为gpu_selector_v）

#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);
}

可以自己尝试修改M，N，K的值（int errCode = gemm(2000, 2000, 2000, 4, 10, my_gpu_queue);中的前三个值），这几个值是修改矩阵的参数，尝试不同的值并比较结果。

下面是M=N=K=2000的示例结果：

实验课练习2：

代码：

（来源：https://github.com/pengzhao-intel/oneAPI_course/blob/main/code/gemm_tile.cpp#L12）（修改同上）

#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）

 

（4，2）

 

 （4，4）

（2，8）

（8，2）

 

（8，8）

 

分析：

1、首先不论x,y的值怎么变化对CPU时间并没有多大的影响（暂且看来，还没有尝试极大的数字）。

2、通过比对（2，2），（2，4），（2，8）这几组发现，y数值变大对于GPU的时间有影响但并不大。

3、通过比对（2，2），（4，2），（8，2）这几组发现，x数值变大对于GPU的时间影响很大。

4、通过比对（2，2），（4，4），（8，8）这几组发现，x,y数值同时变大对GPU时间影响要比单独增大x或y来的大，且有成倍减小趋势，猜测是由于2，4，8都为2的指数倍数。

5、当x，y不是2的整数幂时，没有结果（可自行尝试）。

 

 

破碎的梦782
关注
  • 3
    点赞
  • 踩
  • 4
    收藏
    觉得还不错? 一键收藏
  • 1
    评论
DpC++(DpCpp)入门上手指南
Junital的博客
06-18 673
DPC++是Intel提供的一种编程语言和开发工具，它基于SYCL（异构计算接口标准）并与C++混合使用，旨在简化并行计算在异构系统上的开发。它是oneAPI的一部分，oneAPI是一个开放、跨平台的编程模型，旨在提供统一的编程接口，使开发人员能够在不同的硬件架构上进行高性能计算。使用DPC++，开发人员可以编写能够在多种处理器上运行的代码，包括CPU、GPU、FPGA和其他加速器。
英特尔 oneAPI —— 实现COO格式稀疏矩阵乘积openMP并行化
weixin_50196074的博客
05-17 785
英特尔 oneAPI —— 实现COO格式稀疏矩阵乘积openMP并行化 英特尔 oneAPI 介绍 英特尔 oneAPI 相当于一个超级组件，结合了CPU、GPU、FPGA等等，提供了不同架构和不同硬件的环境，为程序员提供了跨架构的开发能力。 它包括了 oneAPI 标准组件、含有一系列性能库的基于 API 的编程工具，以及先进的分析、调试工具等组件。开发人员可以在英特尔 DevCloud for oneAPI 上对基于多种英特尔架构的代码和应用进行测试。为开发人员的开发提供了相当的便利条件。 提供官网地
DLUT操作系统上机补充实验 —Intel® DevCloud for oneAPI
m0_63707439的博客
06-05 225
Dpcpp的实践环境有三种：1. 基于Intel DevCloud环境2. Dpcpp编译器的本地安装3. 基于Docker的环境这里只介绍前两种环境搭建。
基于DPC++编写矩阵乘法计算代码
m0_63525555的博客
06-01 62
CPU执行时间通过C++标准库中的<chrono>来测量。对代码进行执行时间分析，比较不同实现的效率差异。方法提交一个任务，并在任务中测量事件的执行时间。，并计算二者之间的时间差来得到CPU执行时间。GPU执行时间通过SYCL库中的事件对象的。编写矩阵乘法代码实现，并编译执行。方法获取GPU事件的开始和结束时间。函数中，通过获取开始时间点。类计算得到CPU执行时间。具体是通过调用事件对象的。属性来启用事件的性能分析。类来获取高精度的时间点。，得到GPU执行时间。函数中，通过创建一个。
使用Intel® oneAPI DPC++ 编译器加速矩阵乘法计算
m0_61678439的博客
05-25 638
我们使用SYCL编程模型编写了DPC++代码来实现矩阵乘法，并提供了相应的代码示例。通过使用DPC++编译器，我们可以将矩阵乘法计算加速到在CPU、GPU和FPGA等硬件上运行，这对于大规模数据处理具有很大的意义。首先，我们会简单介绍DPC++编译器的基本概念和使用方法，然后会讲解如何使用DPC++编译器实现矩阵乘法，并提供相应的代码示例。在上述命令中，我们指定了使用DPC++编译器进行编译，使用C++17标准，并指定了使用SYCL编程模型。接下来，我们将编写DPC++代码来实现矩阵乘法。
DLUT操作系统上机补充实验续—矩阵乘法示例
m0_63707439的博客
06-07 106
上一次的实验博客里面写到我的电脑不适配本地的Dpcpp编译器，只能在DevCloud上面写，而DevCloud没有GPU设备，恰巧矩阵乘法示例实验要用到CPU和GPU两个，因此我只能借用舍友电脑来完成实验了。
dlut可视化大作业3————复变函数积分可视化
04-07
【标题】"dlut可视化大作业3————复变函数积分可视化"揭示了本次作业的核心内容，即通过可视化技术展示复变函数的积分过程。在计算机图形学和科学计算领域，可视化是一种强大的工具，它能够帮助我们理解复杂的...
C++PPT（DLUT版含上机作业）
10-28
DLUT版C++上课，内含上机作业！希望能给予您帮助
大连理工大学——J2EE复习资料：复习提纲+题库+ppt
12-01
大连理工J2EE课程，学习资料+复习资料。包括提纲、题库、知识点等。 把手头上的都拖进去了，所以有点乱，请自行选择合适的取用。
基于qt c++Sqlite的 通信录管理系统
01-04
在IT领域，构建一个基于Qt C++和SQLite的通信录管理系统是一项常见的任务，它涉及到用户界面设计、数据库操作以及数据管理等多个方面。Qt是一个跨平台的C++图形用户界面应用程序开发框架，而SQLite则是一个轻量级的...
DlutSpground:这是DLUT Gym的图书系统项目！
06-16
【标题】"DlutSpground:这是DLUT Gym的图书系统项目！" 指的是一个基于Java开发的图书管理系统，由大连理工大学（DLUT）的学生或团队创建。这个项目可能是一个教学实践或者课程设计的成果，目的是为用户提供图书查询...
Simple-Calculator-Android：大连理工大学软件学院20级大一上学期OurEDA实验室安卓组大作业DLUT OurEDA Android分配
02-15
大连理工大学软件学院20级大一上学期OurEDA实验室安卓组大作业 项目说明 实现了一个简易的科学计算器，支持竖版的简易功能和横屏后的科学计算功能 由于只是个作业，以后随缘更新（就是不更新了的意思 另外使用的科学...
计算机网络课件：5_8_5 电子邮件系统.pptx
06-23
电子邮件系统是计算机网络中应用层的一个重要组成部分，它允许用户在不同的计算机和操作系统之间交换信息。在TCP/IP体系中，电子邮件地址由收件人的用户名和邮箱所在主机的域名组成，用"@"符号分隔，例如"user@mail....
计算机网络课件：5_8_4 域名系统DNS.pptx
06-22
【计算机网络】\n\n在计算机网络中，域名系统DNS（Domain Name System）扮演着至关重要的角色。DNS的主要任务是将人类可读的域名转换为网络通信所需的IP地址，使得用户能够方便地访问互联网上的各种服务。本课件将...
英特尔oneAPI—DPCT 移植 CUDA 程序方法、示例及注意事项
qq_15002375的博客
04-23 1938
oneAPI 提供了 DPCT 工具来移植 CUDA 程序，官方称可以移植 95 %的 CUDA 程序。实际体验下来，在自动生成的代码有一些小的语法上错误，需要做一些修正才可正常使用。暂时感觉 AMD 的 hipify 转出来的程序需要修改的内容更少。可以通过以下途径获取和安装 DPCT： 安装 Intel OneAPI Base Toolkit 。 单独安装 DPCT。 DPCT 的使用方式 移植单个 CUDA 文件 单个文件时，对 CUDA 程序进行迁移十分的简单，只需要输入下面的命令行即可完成迁移
操作系统补充实验2 DPC++：矩阵乘法示例
muxuchanyuan的博客
06-17 261
在计算后，通过计时器记录GPU的相乘时间。我们通过对比GEMM_basic 矩阵计算和GEMM_tile 矩阵计算的性能差异，发现在优化后的 GEMM 矩阵计算中，能够充分利用 GPU 的并行计算能力和内存层次结构，大幅提升计算性能。③当tile_X不变时，tile_Y 变化时，当tile_Y修改较小时，如从2->4,4->8,8->16, GPU 的Computation Time也没有很明显的变化，修改逐渐变大时，如16->32,32->64, GPU 的Computation Time会明显增大。
nmf的matlab代码-SDGNMF:此回购实现了JingJing等人的“稀疏双图正则化NMF图像共聚”特征提取的迭代过程。在DLUT
05-26
此代码在人脸上进行聚类实验（ORL-32和PIE-pose27 ）和对象（COIL20）数据集以获得AC和NMI，然后验证SDGNMF的聚类有效性。 此外，该代码还可以显示我们的稀疏性研究和运行时间，以便与其他基于NMF的聚类方法进行比较...
dlut-sql-rar
11-14
大连理工大学提供的这个"dlut-sql-rar"很可能是该大学数据库原理课程的教学资料，包括一些考试的经典试题。这些题目涵盖了数据库的基础知识，如数据管理技术的发展阶段、数据库的三级视图、数据独立性、关系模式、...
dlut数据结构上机
最新发布
10-18
好的，关于 dlut 数据结构上机，我需要更多的信息才能回答你的问题。你可以告诉我具体是哪门课程的上机，或者是哪些数据结构的实现？这样我才能更好地帮助你解答问题。 同时，我需要提醒你，我是一个 AI 对话系统，...

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论 1
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值