跳轉到內容

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://wikibook.tw/w/index.php?title=ROSE_Compiler_Framework/OpenMP_Acclerator_Model_Implementation&oldid=3522602"
華夏公益教科書