SYCL作业四：实验课练习

SWPUMarshall

已于 2023-12-04 20:55:08 修改

阅读量951

点赞数 16

文章标签： oneapi c++

于 2023-12-04 20:54:32 首次发布

本文链接：https://blog.csdn.net/qq_63574067/article/details/134793603

版权

本文介绍了如何在SYCL的gemm_basic代码中，利用workgroup和localworkitem的坐标来计算全局坐标，通过nd_range和nd_item类的函数实现工作组和线程索引的转换。学生需要修改代码以适应M=N=K=2000的输入，并通过正确性测试。

摘要由CSDN通过智能技术生成

SYCL作业四：实验课练习

团队名称

西南石油大学：龙井湖二队

问题陈述

改写gemm_basic代码26，27行，利用work group和local work item的坐标来计算global坐标。

项目源码：https://github.com/pengzhao-intel/oneAPI_course/blob/main/code/gemm_basic.cpp#L26

源码中需要进行修改的部分：

在这里插入图片描述

实现方案

nd_range三维空间示例图：
在这里插入图片描述

本题要求的只是一个二维空间：
在这里插入图片描述

一个橙色部分是一个work-group对应到cuda模型就是一个块，橙色部分的一个小方块是一个item对应到cuda模型就是一个线程。(我们当前举例的是work-group的大小是4*4而不是必须是4*4)

所以本题考察怎样通过work-group和item的索引来确定一个item在全局的索引。

比如我要确定这个点的坐标
在这里插入图片描述
那么就要看他是在第几个work-group，和他在当前work-group中第几个item。由图显然他在第3个work-group中行为1，列为1（work-group由0开始计数）。所以就可以由此得出他在整体的索引

rows = 1 * (4*4)+ 1;

cols = 1 * (4*4) + 1;

对应到cuda当中就是

rows = blockIdx.x * blockDim.x + threadIdx.x；

cols = blockIdx.y * blockDim.y + threadIdx.y;

那么oneapi里面哪些参数对应一个work-group的大小，该item在哪个work-group中，该item在一个work-group中的x方向的索引，y方向的索引呢。又该怎么用呢。

在oneapi中
在这里插入图片描述
创建一个二维的nd模型x轴对应global_ndrange,y轴对应local_ndrange,并用nd_item类来映射二维模型

nd_item类中又有很多方法

template <int Dimensions = 1> class nd_item {
 public:
  static constexpr int dimensions = Dimensions;

  nd_item() = delete;

  /* -- common interface members -- */

  id<Dimensions> get_global_id() const;

  size_t get_global_id(int dimension) const;

  size_t get_global_linear_id() const;

  id<Dimensions> get_local_id() const;

  size_t get_local_id(int dimension) const;

  size_t get_local_linear_id() const;

  group<Dimensions> get_group() const;

  sub_group get_sub_group() const;

  size_t get_group(int dimension) const;

  size_t get_group_linear_id() const;

  range<Dimensions> get_group_range() const;

  size_t get_group_range(int dimension) const;

  range<Dimensions> get_global_range() const;

  size_t get_global_range(int dimension) const;

  range<Dimensions> get_local_range() const;

  size_t get_local_range(int dimension) const;

  // Deprecated in SYCL 2020.
  id<Dimensions> get_offset() const;

  nd_range<Dimensions> get_nd_range() const;
};
}

网上查询每个函数得知：

get_local_id(int dimension);//返回该item所在work-group中dimension方向上的索引

get_group(int dimension);//返回该item是在dimension方向上哪个work-group

get_local_range(int dimension);//返回一个work-group中dimension方向上有多少个item

便由此推理出以下结果

在这里插入图片描述