跳转到内容

ROSE 编译器框架/OpenMP 加速器模型实现

Add links
来自维基教科书，开放的世界，开放的书籍

ROSE 编译器框架/OpenMP 支持>

概述

[编辑 | 编辑源代码]

我们正在试验 OpenMP 4.0 规范中 OpenMP 加速器扩展的试用实现。由于规范快速变化的性质和我们的资源限制，它并非官方或完整的实现。

该实现基于 ROSE 中现有的 OpenMP 实现。

已实现的功能

[编辑 | 编辑源代码]

该实现大体遵循技术报告 http://www.openmp.org/mp-documents/TR1_167.pdf 和后来的 http://www.openmp.org/mp-documents/OpenMP4.0.0.pdf 。

  • pragma omp target + pragma omp parallel for
  • map 子句
  • GPU 上的 reduction
  • pragma omp target data
  • collapse 子句

编译和运行

[编辑 | 编辑源代码]

我们逐步发布这项工作，作为托管在 https://github.com/rose-compiler/rose-develop 的 ROSE（基于 EDG 4.x 的版本）编译器的一部分。

  • 克隆它
  • 在 buildtree/src 下配置和构建 librose.so
  • 通过键入 "make roseompacc" 构建翻译器 roseompacc，位于 tests/nonsmoke/functional/roseTests/ompLoweringTests/

翻译器源代码位于 rose-omp-accelerator/tests/nonsmoke/functional/roseTests/ompLoweringTests/roseompacc.C

使用 roseompacc 的命令行（目前分为两步）在 rose-omp-accelerator/tests/nonsmoke/functional/roseTests/ompLoweringTests/Makefile.am 中举例说明。

# Experimental translation for OpenMP accelerator model directives
# no final compilation for now, which requires CUDA compiler for the generated code
test_omp_acc:axpy_ompacc.o matrixmultiply-ompacc.o jacobi-ompacc.o
rose_axpy_ompacc.cu:roseompacc
        ./roseompacc$(EXEEXT) ${TEST_FLAGS} -rose:openmp:lowering -rose:skipfinalCompileStep -c $(TEST_DIR)/axpy_ompacc.c 
rose_matrixmultiply-ompacc.cu:roseompacc
        ./roseompacc$(EXEEXT) ${TEST_FLAGS} -rose:openmp:lowering -rose:skipfinalCompileStep -c $(TEST_DIR)/matrixmultiply-ompacc.c 
rose_jacobi-ompacc.cu:roseompacc
        ./roseompacc$(EXEEXT) ${TEST_FLAGS} -rose:openmp:lowering -rose:skipfinalCompileStep -c $(TEST_DIR)/jacobi-ompacc.c 

# build executables using nvcc
axpy_ompacc.out:rose_axpy_ompacc.cu
        nvcc $<  $(TEST_INCLUDES) $(top_srcdir)/src/midend/programTransformation/ompLowering/xomp_cuda_lib.cu -o $@
matrixmultiply-ompacc.out:rose_matrixmultiply-ompacc.cu
        nvcc $< $(TEST_INCLUDES) $(top_srcdir)/src/midend/programTransformation/ompLowering/xomp_cuda_lib.cu -o $@
jacobi-ompacc.out:rose_jacobi-ompacc.cu
        nvcc $< $(TEST_INCLUDES) $(top_srcdir)/src/midend/programTransformation/ompLowering/xomp_cuda_lib.cu -o $@

实现细节

[编辑 | 编辑源代码]

解析

[编辑 | 编辑源代码]

flex/bison

存储解析结果的 AST 属性

IR 生成：ROSETTA，将属性转换为专用 AST

翻译

[编辑 | 编辑源代码]

翻译

在内部，它调用以经典模式（传递单个参数，没有任何包装数组或结构）运行的轮廓器

  • Outliner::enable_classic = true;

运行时支持

[编辑 | 编辑源代码]

运行时支持

关于将什么放入每个文件的规则

  • libxomp.h 编译器编写者的运行时接口
  • xomp.c CPU 线程的运行时支持
  • xomp_cuda_lib.cu 提供了由主机代码调用的包装 CUDA 代码（没有 __device__ 关键字）的常规 C 函数。
  • xomp_cuda_lib_inlined.cu 提供了由 CUDA 内核函数（具有 __global__ 关键字）调用的 __device__ 函数。

在进行这项工作时，CUDA 和/或 nvcc 没有针对设备代码的链接器（链接来自不同文件的 __global__ 和 _device_ 函数）。我们必须将一些常见的设备函数放入此文件中。这样，它们就可以包含在生成 CUDA 内核的同一个文件中（__global__ 函数只能调用来自同一个文件的 __device__ 函数）。

设备数据管理

[编辑 | 编辑源代码]

运行时提供两种管理设备数据的方法

  • 显式管理数据分配、复制到、复制回和释放
  • 使用设备数据环境 (DDE) 函数自动管理
    • xomp_deviceDataEnvironmentEnter(); // 进入环境，推送到环境堆栈
    • xomp_deviceDataEnvironmentPrepareVariable(); // 如果已在父环境中，则重用数据，或分配，并向当前环境堆栈注册
    • xomp_deviceDataEnvironmentExit(); // 自动释放，复制回数据。弹出环境堆栈

显式数据管理示例

int mmm()
{
{
    float *_dev_a;
    int _dev_a_size = sizeof(float ) * N * M;
    _dev_a = ((float *)(xomp_deviceMalloc(_dev_a_size)));
    xomp_memcpyHostToDevice(((void *)_dev_a),((const void *)a),_dev_a_size);
    float *_dev_b;
    int _dev_b_size = sizeof(float ) * M * K;
    _dev_b = ((float *)(xomp_deviceMalloc(_dev_b_size)));
    xomp_memcpyHostToDevice(((void *)_dev_b),((const void *)b),_dev_b_size);
    float *_dev_c;
    int _dev_c_size = sizeof(float ) * N * M;
    _dev_c = ((float *)(xomp_deviceMalloc(_dev_c_size)));
    xomp_memcpyHostToDevice(((void *)_dev_c),((const void *)c),_dev_c_size);
/* Launch CUDA kernel ... */
    int _threads_per_block_ = xomp_get_maxThreadsPerBlock();
    int _num_blocks_ = xomp_get_max1DBlock(1023 - 0 + 1);
    OUT__1__9221__<<<_num_blocks_,_threads_per_block_>>>(_dev_a,_dev_b,_dev_c);
    xomp_freeDevice(_dev_a);
    xomp_freeDevice(_dev_b);
    xomp_memcpyDeviceToHost(((void *)c),((const void *)_dev_c),_dev_c_size);
    xomp_freeDevice(_dev_c);
  }
  return 0;
}

自动数据管理支持重用嵌套设备数据环境

int mmm()
{
{
    xomp_deviceDataEnvironmentEnter();
    float *_dev_a;
    int _dev_a_size = sizeof(float ) * 1024 * 1024;
    _dev_a = ((float *)(xomp_deviceDataEnvironmentPrepareVariable(((void *)a),_dev_a_size,1,0)));
    float *_dev_b;
    int _dev_b_size = sizeof(float ) * 1024 * 1024;
    _dev_b = ((float *)(xomp_deviceDataEnvironmentPrepareVariable(((void *)b),_dev_b_size,1,0)));
    float *_dev_c;
    int _dev_c_size = sizeof(float ) * 1024 * 1024;
    _dev_c = ((float *)(xomp_deviceDataEnvironmentPrepareVariable(((void *)c),_dev_c_size,1,1)));
/* Launch CUDA kernel ... */
    int _threads_per_block_ = xomp_get_maxThreadsPerBlock();
    int _num_blocks_ = xomp_get_max1DBlock(1023 - 0 + 1);
    OUT__1__9221__<<<_num_blocks_,_threads_per_block_>>>(_dev_a,_dev_b,_dev_c);
    xomp_deviceDataEnvironmentExit();
  }
  return 0;
}

环境变量

[编辑 | 编辑源代码]

默认情况下控制在多 GPU 支持中使用的设备数量

 export OMP_NUM_DEVICES=5

运行时会自动检测此环境变量并使用它来控制 GPU 计数。

你可以在代码中检索此计数，例如

 int GPU_N = xomp_get_num_devices(); // this function will obtain the env variable.

如果指定了 OMP_NUM_DEVICES，xomp_get_num_devices() 将在内部调用 int omp_get_max_devices(void) 来获取硬件支持的最大设备数量。

相关 提交

循环调度

[编辑 | 编辑源代码]

必须提供运行时循环调度，因为 CUDA 教程中简单的 1-迭代到 1-线程映射对于大型迭代空间将不起作用。（迭代计数 > 总 GPU 线程）。

// Jacobi 2-D kernel: 
// Naive 1-iteration to 1 GPU thread scheduling: each GPU thread gets one iteration to work on.
// Problem: won't scale to large iteration space.
    int i = blockIdx.x * blockDim.x + threadIdx.x + 1;
    int j = blockIdx.y * blockDim.y + threadIdx.y + 1;

    newa[j*m+i] = w0*a[j*m+i] +
	    w1 * (a[j*m+i-1] + a[(j-1)*m+i] +
		  a[j*m+i+1] + a[(j+1)*m+i]) +
	    w2 * (a[(j-1)*m+i-1] + a[(j+1)*m+i-1] +
		  a[(j-1)*m+i+1] + a[(j+1)*m+i+1]);


我们在运行时提供了两个循环调度器，因此每个 CUDA 线程都可以处理多个迭代来处理大型迭代空间。

  1. 静态偶数调度器：此调度器将迭代空间均匀地划分为大致相等大小的块。然后，它将每个块分配给一个 CUDA 线程。这种调度策略可能会过度使用内存，因为每个线程将接触由迭代块带来的很大范围的数据。
  2. 循环调度器：每个线程一次获取一个迭代（或多个迭代）。它与 OpenMP 的 schedule(static, chunk) 策略相同。

调度器的测试使用（使用相应的 CPU 版本）在以下位置给出：

GPU 版本的使用示例：循环调度器（现在是新的默认调度器）

void OUT__2__10550__(int n,int *_dev_u)
{
  int ij;
  int _dev_lower, _dev_upper;


  // variables for adjusted loop info considering both original chunk size and step(strip)
  int _dev_loop_chunk_size;
  int _dev_loop_sched_index;
  int _dev_loop_stride;

  // 1-D thread block:
  int _dev_thread_num = omp_get_num_threads();
  int _dev_thread_id = omp_get_thread_num();
  printf ("thread count = %d, current thread id = %d\n", _dev_thread_num, _dev_thread_id);

  int orig_start =0; // must be correct!!
  int orig_end = n-1; // use inclusive bound
  int orig_step = 1;
  int orig_chunk_size = 1;

  XOMP_static_sched_init (orig_start, orig_end, orig_step, orig_chunk_size, _dev_thread_num, _dev_thread_id, \
      & _dev_loop_chunk_size , & _dev_loop_sched_index, & _dev_loop_stride);

  printf ("Initialized chunk size = %d, sched indx =%d, stride = %d\n",_dev_loop_chunk_size, _dev_loop_sched_index, _dev_loop_stride);

  while (XOMP_static_sched_next (&_dev_loop_sched_index, orig_end, orig_step, _dev_loop_stride, _dev_loop_chunk_size, _dev_thread_num, _dev_thread_id, & _dev_lower
        , & _dev_upper))
  {
    printf ("Thread ID: %d Allocated lower = %d upper = %d\n", _dev_thread_id, _dev_lower, _dev_upper);
    for (ij = _dev_lower ; ij <= _dev_upper; ij ++) { // using inclusive bound here
        _dev_u[ij] += (n - ij);         
    }
  }
}

如何使用静态-偶调度器的示例（由于性能缺陷，在最近的版本中不再是默认调度器。）

// using a scheduler
__global__ void OUT__1__11058__(int j,int k,float *_dev_a,float *_dev_b,float *_dev_c)
{
  int _dev_i;
  long _dev_lower, _dev_upper;
  XOMP_accelerator_loop_default (0, MSIZE -1 , 1, &_dev_lower, &_dev_upper);

  for (_dev_i = _dev_lower; _dev_i<= _dev_upper; _dev_i ++) 
  {
    for (j = 0; j < MSIZE; j++)
    {
      float c= 0.0;
      for (k = 0; k < MSIZE; k++)
        c += _dev_a[_dev_i * MSIZE + k] * _dev_b[k * MSIZE + j];
      _dev_c[_dev_i * MSIZE + j] = c;
    }
  }
}

测试

[编辑 | 编辑源代码]

输入文件位于 https://github.com/rose-compiler/rose-develop/tree/master/tests/nonsmoke/functional/CompileTests/OpenMP_tests

  • axpy_ompacc.c
  • axpy_ompacc2.c
  • matrixmultiply-ompacc.c
  • jacobi-ompacc.c
  • jacobi-ompacc-opt1.c // 使用“目标数据”区域

测试目录 https://github.com/rose-compiler/rose-develop/tree/master/tests/nonsmoke/functional/roseTests/ompLoweringTests

  • 测试翻译器：roseompacc.C
  • 测试驱动程序：Makefile.am
  • 在构建树中执行 make ompacc_test 将触发相关测试
    • 确保您已安装 nvidia nvcc 编译器以生成可执行文件。否则，只会生成 cuda 文件，最终编译将失败。

一些示例生成的 CUDA 文件可以在以下位置找到：

出版物

[编辑 | 编辑源代码]
检索自 "https://wikibooks.cn/w/index.php?title=ROSE_Compiler_Framework/OpenMP_Acclerator_Model_Implementation&oldid=3522602"
华夏公益教科书