スポンサーサイト

上記の広告は1ヶ月以上更新のないブログに表示されています。
新しい記事を書く事で広告が消せます。

改めて、CUDA BY EXAMPLEのサンプルをOpenCLで

標題の通り、CUDA BY EXAMPLEという書籍がある。原著はAddison-Wesley、日本語版はインプレスジャパンから出ている。NVIDIAの技術者によって書かれているだけあって、薄い本ながらいかにして並列処理の性能を引き出すかについて注意が払われている。

私の今使用している環境は主にLinux + OpenCLなので、この本のサンプルのうち何度も取り上げられているベクトルの内積の例を自分の環境に合わせて翻案し、テストしてみたい。

ということで、まずは原型として、chapter05/dot.cuに相当するプログラムを作成する。CUDA Cではホスト(CPU)側コードもデバイス(GPU)側コードも同一のソースに記述することになるが、OpenCLではGLSLのようなシェーダー言語を使う場合に似て、両者を別々のファイルで管理することになる。ついでに言えば、OpenCLのAPIはデバイス側のリソース管理は基本的にOpenGLでのテクスチャオブジェクトやバッファオブジェクト等とよく似た形式で管理されるので、OpenGLに馴染みがある人はOpenCLのAPIにもすぐ親しみがわくであろう。

というわけで、まずホスト側コード(main.cpp)

#include <iostream>
#include <CL/opencl.h>
#include "myCLUtils.h"
#include "myCLManager2.h"

#define N (33 * 1024)
#define THREADSPERBLOCK 256
#define imin(a, b) (a < b ? a : b)
#define expectedSum(x) (x * (x + 1)*(2 * x + 1) / 3)

int main(int argc, char **argv)
{
  cl_int ret;
  int i, n = N;
  const int queueNum = 1, debug = 0;
  const int threadsPerBlock = THREADSPERBLOCK;
  const int blocksPerGrid
    = imin(32, (N + THREADSPERBLOCK - 1) / THREADSPERBLOCK);
  const size_t globalWorkSize[1] = {N};
  const size_t localWorkSize[1] = {threadsPerBlock};

  float *a, *b, *partial_c;
  cl_mem dev_a, dev_b, dev_partial_c;
  myCLManager2 clm(queueNum, debug);
  clm.register_program("dotSource", "dotProduct.cl");
  cl_program program = clm.get_program("dotSource");
  cl_kernel dotKernel = clCreateKernel(program, "dotProduct", &ret);
  myCLUtils::report_error(__FUNCTION__, "clCreateKernel", ret);
  a = new float[N];
  b = new float[N];
  partial_c = new float[blocksPerGrid];
  for(i = 0; i < N; i++){
    a[i] = i;
    b[i] = i * 2;
  }
  dev_a = clCreateBuffer(clm.context,
             CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
             N * sizeof(float), a, &ret);
  myCLUtils::report_error(__FUNCTION__, "clCreateBuffer", ret);
  dev_b = clCreateBuffer(clm.context,
             CL_MEM_READ_ONLY|CL_MEM_COPY_HOST_PTR,
             N * sizeof(float), b, &ret);
  myCLUtils::report_error(__FUNCTION__, "clCreateBuffer", ret);
  dev_partial_c = clCreateBuffer(clm.context,
                 CL_MEM_WRITE_ONLY,
                 blocksPerGrid * sizeof(float), NULL, &ret);
  myCLUtils::report_error(__FUNCTION__, "clCreateBuffer", ret);
  i = 0;
  ret = clSetKernelArg(dotKernel, i++, sizeof(dev_a), &dev_a);
  myCLUtils::report_error(__FUNCTION__, "clSetKernelArg", ret);
  ret = clSetKernelArg(dotKernel, i++, sizeof(dev_b), &dev_b);
  myCLUtils::report_error(__FUNCTION__, "clSetKernelArg", ret);
  ret = clSetKernelArg(dotKernel, i++, sizeof(dev_partial_c), &dev_partial_c);
  myCLUtils::report_error(__FUNCTION__, "clSetKernelArg", ret);
  ret = clSetKernelArg(dotKernel, i++, sizeof(n), &n);
  myCLUtils::report_error(__FUNCTION__, "clSetKernelArg", ret);
  ret = clSetKernelArg(dotKernel, i++,
               sizeof(threadsPerBlock), &threadsPerBlock);
  myCLUtils::report_error(__FUNCTION__, "clSetKernelArg", ret);
  ret = clSetKernelArg(dotKernel, i++,
               sizeof(blocksPerGrid), &blocksPerGrid);
  myCLUtils::report_error(__FUNCTION__, "clSetKernelArg", ret);
  ret = clEnqueueNDRangeKernel(clm.queues[0],
                   dotKernel,
                   1,
                   NULL,
                   globalWorkSize,
                   localWorkSize,
                   0, NULL, NULL);
  myCLUtils::report_error(__FUNCTION__, "clEnqueueNDRangeKernel", ret);
  ret = clEnqueueReadBuffer(clm.queues[0],
                dev_partial_c,
                CL_TRUE,
                0,
                blocksPerGrid * sizeof(float),
                partial_c,
                0, NULL, NULL);
  myCLUtils::report_error(__FUNCTION__, "clEnqueueReadBuffer", ret);
  float c = 0;
  for(int i = 0; i < blocksPerGrid; i++){
    c += partial_c[i];
  }
  delete a;
  delete b;
  delete partial_c;
  ret = clReleaseMemObject(dev_a);
  myCLUtils::report_error(__FUNCTION__, "clReleaseMemObject", ret);
  clReleaseMemObject(dev_b);
  myCLUtils::report_error(__FUNCTION__, "clReleaseMemObject", ret);
  clReleaseMemObject(dev_partial_c);
  myCLUtils::report_error(__FUNCTION__, "clReleaseMemObject", ret);
  cout << "Blocks per Grid = " << blocksPerGrid << "\n";
  cout << "Threads per Block = " << threadsPerBlock << "\n";
  cout << "GPU sum  = " << c << "\n";
  cout << "Expected = " << expectedSum((float)(N - 1)) << endl;
  return 0;
}

そして、GPU側コード(dotProduct.cl)

#pragma OPENCL EXTENSION cl_khr_fp64: enable
#define THREADSPERBLOCK 256

__kernel void dotProduct(__global const float* a,
             __global const float* b,
             __global float *partial_c,
             int n,
             int threadsPerBlock,
             int blocksPerGrid)
{
  int i;
  local float cache[THREADSPERBLOCK];
  int tid = get_global_id(0);
  int cacheIndex = get_local_id(0);
  int gid = get_group_id(0);
  float temp = 0;

  while(tid < n){
    temp += a[tid] * b[tid];
    tid += threadsPerBlock * blocksPerGrid;
  }

  cache[cacheIndex] = temp;
  barrier(CLK_LOCAL_MEM_FENCE);

  i = threadsPerBlock / 2;

  while( i != 0 ){
    if(cacheIndex < i)
      cache[cacheIndex] += cache[cacheIndex + i];
    barrier(CLK_LOCAL_MEM_FENCE);
    i /= 2;
  }

  if(cacheIndex == 0){
    partial_c[gid] = cache[0];
  }
}

あまり高度なことはしていない。エラー処理関数はmyCLUtils名前空間で定義してある。ヘルパークラスとしてmyCLDevice、myCLPlatform、myCLManager2の3つを作成、使用している。myCLDeviceとmyCLPlatformは単にそれぞれデバイスとプラットフォームに関して取得した情報のostreamへの出力を定義しているだけである。
myCLManager2では、内部でOpenCLコンテキストを作成し、それに対して、コンストラクタ引数で指定された個数のコマンドキューを作成するようにしている。

myCLManageer2::MyCLManager2(const int numQueues = 1, const int verbose = 0);

というシグネチャになっていて、第一引数がキューの数、第二引数は単にデバッグのための情報出力フラグである。なお、このサンプルではまだキューは一つだけしか使わない。

特に他人が動かすことを想定しているわけではないが、一応全体のソースはopaquelight.com/MultiQueue1.tgzに置いておくことにする。出力はおおむねこのようになるはずである。

$ ./multiqueue1
Blocks per Grid = 32
Threads per Block = 256
GPU sum  = 2.57236e+13
Expected = 2.57236e+13

これだけでは何の意味もないが、この後上述の書籍の例に沿ってこのサンプルを変更していき、パフォーマンスの検討などをしてみることにしよう。
スポンサーサイト

テーマ : プログラミング
ジャンル : title="コンピュータ">コンピュータ

コメントの投稿

非公開コメント

プロフィール

Author:GM3D
FC2ブログへようこそ!

最新記事
最新コメント
最新トラックバック
月別アーカイブ
カテゴリ
FC2カウンター
検索フォーム
RSSリンクの表示
リンク
ブロとも申請フォーム

この人とブロともになる

QRコード
QR
上記広告は1ヶ月以上更新のないブログに表示されています。新しい記事を書くことで広告を消せます。