質問をすることでしか得られない、回答やアドバイスがある。

15分調べてもわからないことは、質問しよう!

ただいまの
回答率

90.33%

  • C

    3991questions

    C言語は、1972年にAT&Tベル研究所の、デニス・リッチーが主体となって作成したプログラミング言語です。 B言語の後継言語として開発されたことからC言語と命名。そのため、表記法などはB言語やALGOLに近いとされています。 Cの拡張版であるC++言語とともに、現在世界中でもっとも普及されているプログラミング言語です。

  • C++

    3768questions

    C++はC言語をもとにしてつくられた最もよく使われるマルチパラダイムプログラミング言語の1つです。オブジェクト指向、ジェネリック、命令型など広く対応しており、多目的に使用されています。

  • MacOS(OSX)

    2037questions

    MacOSとは、Appleの開発していたGUI(グラフィカルユーザーインターフェース)を採用したオペレーションシステム(OS)です。Macintoshと共に、市場に出てGUIの普及に大きく貢献しました。

  • 並列処理

    39questions

    複数の計算が同時に実行される手法

OpenCLでの二次元配列の処理の仕方について

受付中

回答 1

投稿 編集

  • 評価
  • クリップ 0
  • VIEW 796

khsh

score 2

前提・実現したいこと

OpenCLで+1するプログラムを作成してみました。
https://peta.okechan.net/blog/archives/2538
を参考にしています。

int index = get_global_id(0);
int index2 = get_global_id(1);

を用いて
data[index*M+index2] += 1.0;
のようにアクセスしたいと思っています。

発生している問題・エラーメッセージ

int index = get_global_id(0);
int index2 = get_global_id(1);
を行うと常にindexとindex2に同じ値が格納されてしまいます。
ほとんどの部分が+1されず終わっています。

###該当のソースコード
//
//  main3.cpp
//
#include <iostream>
#include <vector>
#include <OpenCL/opencl.h>
#include <numeric>

#define PLATFORM_MAX 4
#define DEVICE_MAX 4


void EC(cl_int result, const char *title)
{
    if (result != CL_SUCCESS) {
        std::cout << "Error: " << title << "(" << result << ")\n";
    }
}


cl_int err = CL_SUCCESS;
void EC2(const char *title)
{
    if (err != CL_SUCCESS) {
        std::cout << "Error: " << title << "(" << err << ")\n";
    }
    err = CL_SUCCESS;
}


int main(int argc, const char * argv[])
{
    // プラットフォーム一覧を取得
    cl_platform_id platforms[PLATFORM_MAX];
    cl_uint platformCount;
    EC(clGetPlatformIDs(PLATFORM_MAX, platforms, &platformCount), "clGetPlatformIDs");
    if (platformCount == 0) {
        std::cerr << "No platform.\n";
        return EXIT_FAILURE;
    }

    // 見つかったプラットフォームの情報を印字
    for (int i = 0; i < platformCount; i++) {
        char vendor[100] = {0};
        char version[100] = {0};
        EC(clGetPlatformInfo(platforms[i], CL_PLATFORM_VENDOR, sizeof(vendor), vendor, nullptr), "clGetPlatformInfo");
        EC(clGetPlatformInfo(platforms[i], CL_PLATFORM_VERSION, sizeof(version), version, nullptr), "clGetPlatformInfo");
        std::cout << "Platform id: " << platforms[i] << ", Vendor: " << vendor << ", Version: " << version << "\n";
    }

    // デバイス一覧を取得
    cl_device_id devices[DEVICE_MAX];
    cl_uint deviceCount;
    EC(clGetDeviceIDs(platforms[0], CL_DEVICE_TYPE_GPU, DEVICE_MAX, devices, &deviceCount), "clGetDeviceIDs");
    if (deviceCount == 0) {
        std::cerr << "No device.\n";
        return EXIT_FAILURE;
    }

    // 見つかったデバイスの情報を印字
    std::cout << deviceCount << " device(s) found.\n";
    for (int i = 0; i < deviceCount; i++) {
        char name[100] = {0};
        size_t len;
        EC(clGetDeviceInfo(devices[i], CL_DEVICE_NAME, sizeof(name), name, &len), "clGetDeviceInfo");
        std::cout << "Device id: " << i << ", Name: " << name << "\n";
    }

    // コンテキストの作成
    cl_context ctx = clCreateContext(nullptr, 1, devices, nullptr, nullptr, &err);
    EC2("clCreateContext");

    // コンパイル済みclプログラムの読み込み
    const char* bitcode_path = "kernel2.cl.gpu_32.bc";
    size_t len = strlen(bitcode_path);
    cl_program program = clCreateProgramWithBinary(ctx, 1, devices, &len, (const unsigned char**)&bitcode_path, nullptr, &err);
    EC2("clCreateProgramWithBinary");

    // プログラムのビルド
    EC(clBuildProgram(program, 1, devices, nullptr, nullptr, nullptr), "clBuildProgram");

    // カーネルの作成
    cl_kernel kernel = clCreateKernel(program, "addone", &err);
    EC2("clCreateKernel");

    // データを用意
    int n = 10;
    std::vector<float> data(n*n,0.0f);
    // デバイスメモリを確保しつつデータをコピー
    cl_mem device_mem = clCreateBuffer(ctx, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(float) * n*n, data.data(), &err);
    EC2("clCreateBuffer");

    // カーネルの引数をセット
    EC(clSetKernelArg(kernel, 0, sizeof(cl_mem), &device_mem), "clSetKernelArg");
    EC(clSetKernelArg(kernel, 1, sizeof(int), &n), "clSetKernelArg");

    // コマンドキューの作成
    cl_command_queue q = clCreateCommandQueue(ctx, devices[0], 0, &err);
    EC2("clCreateCommandQueue");

    // カーネルの実行
    size_t global[2],local[2],offset[2];
    offset[0] = 0;
    offset[0] = 0;
    global[0] = n;
    global[1] = n;
    local[0] = 1;
    local[1] = 1;
    EC(clEnqueueNDRangeKernel(q, kernel, 2, offset, global, NULL, 0, nullptr, nullptr), "clEnqueueNDRangeKernel");

    // 結果を読み込み
    EC(clEnqueueReadBuffer(q, device_mem, CL_TRUE, 0, sizeof(float) * n*n, data.data(), 0, nullptr, nullptr), "clEnqueueReadBuffer");

    // 結果の印字
    for (int i = 0; i < n*n; i++) {
        std::cout << data[i] << ", ";
    }
    std::cout << "\n";
    float total = std::accumulate(data.begin(),data.end(),0.0);
    std::cout << total << std::endl;
    // コマンドキューの解放
    EC(clReleaseCommandQueue(q), "clReleaseCommandQueue");

    // デバイスメモリを解放
    EC(clReleaseMemObject(device_mem), "clReleaseMemObject");

    // カーネルの解放
    EC(clReleaseKernel(kernel), "clReleaseKernel");

    // プログラムの解放
    EC(clReleaseProgram(program), "clReleaseProgram");

    // コンテキストの解放
    EC(clReleaseContext(ctx), "clReleaseContext");

    std::cout << "Done.\n";
    return EXIT_SUCCESS;
}



//ここからカーネル部分
//kernel2.cl
__kernel
void addone(__global float* data,const int n)
{
    int index = get_global_id(0);
    int index2 = get_global_id(1);
    int dim = get_work_dim();
    printf("get_work_dim = %d\n",dim);
    printf("index = %d , index2 = %d \n",index,index2);
    data[index*n+index2] += 1.0f;
}

C++
```

試したこと

EC(clEnqueueNDRangeKernel(q, kernel, 2, offset, global, local, 0, nullptr, nullptr), "clEnqueueNDRangeKernel");のところのglobal,localなどの数値を変えるなどして動かしてみました。

補足情報(言語/FW/ツール等のバージョンなど)

MacBook Pro 13インチ、OSはel capitanです。

Terminal上で
/System/Library/Frameworks/OpenCL.framework/Libraries/openclc -c -o kernel2.cl.gpu_32.bc -arch gpu_32 -emit-llvm kernel2.cl
g++ -O3 -std=c++11 -framework opencl main3.cpp -o test
./test
で実行しました。

  • 気になる質問をクリップする

    クリップした質問は、後からいつでもマイページで確認できます。

    またクリップした質問に回答があった際、通知やメールを受け取ることができます。

    クリップを取り消します

  • 良い質問の評価を上げる

    以下のような質問は評価を上げましょう

    • 質問内容が明確
    • 自分も答えを知りたい
    • 質問者以外のユーザにも役立つ

    評価が高い質問は、TOPページの「注目」タブのフィードに表示されやすくなります。

    質問の評価を上げたことを取り消します

  • 評価を下げられる数の上限に達しました

    評価を下げることができません

    • 1日5回まで評価を下げられます
    • 1日に1ユーザに対して2回まで評価を下げられます

    質問の評価を下げる

    teratailでは下記のような質問を「具体的に困っていることがない質問」、「サイトポリシーに違反する質問」と定義し、推奨していません。

    • プログラミングに関係のない質問
    • やってほしいことだけを記載した丸投げの質問
    • 問題・課題が含まれていない質問
    • 意図的に内容が抹消された質問
    • 広告と受け取られるような投稿

    評価が下がると、TOPページの「アクティブ」「注目」タブのフィードに表示されにくくなります。

    質問の評価を下げたことを取り消します

    この機能は開放されていません

    評価を下げる条件を満たしてません

    評価を下げる理由を選択してください

    詳細な説明はこちら

    上記に当てはまらず、質問内容が明確になっていない質問には「情報の追加・修正依頼」機能からコメントをしてください。

    質問の評価を下げる機能の利用条件

    この機能を利用するためには、以下の事項を行う必要があります。

回答 1

0

//ここからカーネル部分 
//kernel2.cl 
kernel 
 void addone(global float* data,const int n) { 
   int index = get_global_id(1); // ← ココ、0じゃないの?
   int index2 = get_global_id(1); 
   printf("index = %d , index2 = %d \n",index,index2); 
   data[index*n+index2] += 1.0f; 
 } 

投稿

  • 回答の評価を上げる

    以下のような回答は評価を上げましょう

    • 正しい回答
    • わかりやすい回答
    • ためになる回答

    評価が高い回答ほどページの上位に表示されます。

  • 回答の評価を下げる

    下記のような回答は推奨されていません。

    • 間違っている回答
    • 質問の回答になっていない投稿
    • スパムや攻撃的な表現を用いた投稿

    評価を下げる際はその理由を明確に伝え、適切な回答に修正してもらいましょう。

  • 2017/01/11 10:32

    申しわけありません。カーネルの部分を色々試しており、修正するのを忘れていました。
    int index = get_global_id(1) を
    int index = get_global_id(0) に変えても問題は解決しませんでした。

    キャンセル

  • 2017/01/11 10:35

    kernel内の printf はどんな出力するんですか? 10*10 = 100行出力されますか?

    キャンセル

  • 2017/01/11 10:46

    もうひとつ。
    EC(clSetKernelArg(kernel, 1, sizeof(cl_mem), &n), "clSetKernelArg");
                   ↑sizeof(int) じゃないの?

    キャンセル

  • 2017/01/11 11:08

    index = 0 , index2 = 0
    index = 0 , index2 = 0
    index = 0 , index2 = 0
    index = 1 , index2 = 1
    index = 0 , index2 = 0
    index = 0 , index2 = 0
    index = 1 , index2 = 1
    index = 2 , index2 = 2
    途中省略
    index = 9 , index2 = 9
    index = 8 , index2 = 8
    index = 9 , index2 = 9
    index = 9 , index2 = 9
    index = 9 , index2 = 9
    index = 8 , index2 = 8
    のような出力で100行出力されていました。

    EC(clSetKernelArg(kernel, 1, sizeof(cl_mem), &n), "clSetKernelArg")を
    EC(clSetKernelArg(kernel, 1, sizeof(int), &n), "clSetKernelArg")に変えましたが結果は変わりませんでした。

    キャンセル

  • 2017/01/11 11:21 編集

    ...わからんなー、100行書かれたんなら n*n 個のkernelが動いたわけで、なんで get_global_id() が同じ値になるんやろ。get_work_dim() は2を返してくれますか?

    # ごめんなさい、一昨年OpenCLからCUDAに寝返ったので記憶が曖昧。これ以上はアドバイスできんです。

    キャンセル

  • 2017/01/11 12:22

    kernel2.cl上でget_work_dim()をしてみたところget_work_dim() は2を返していました。

    アドバイスありがとうございました。

    キャンセル

  • 2017/01/11 12:28

    ダメ元で clEnqueueNDRangeKernelの第6パラメータ: local を NULL にすれば?

    キャンセル

  • 2017/01/11 12:43

    clEnqueueNDRangeKernelの第6パラメータをNULLにしてみましたが、結果は変わりませんでした。

    キャンセル

15分調べてもわからないことは、teratailで質問しよう!

  • ただいまの回答率 90.33%
  • 質問をまとめることで、思考を整理して素早く解決
  • テンプレート機能で、簡単に質問をまとめられる

同じタグがついた質問を見る

  • C

    3991questions

    C言語は、1972年にAT&Tベル研究所の、デニス・リッチーが主体となって作成したプログラミング言語です。 B言語の後継言語として開発されたことからC言語と命名。そのため、表記法などはB言語やALGOLに近いとされています。 Cの拡張版であるC++言語とともに、現在世界中でもっとも普及されているプログラミング言語です。

  • C++

    3768questions

    C++はC言語をもとにしてつくられた最もよく使われるマルチパラダイムプログラミング言語の1つです。オブジェクト指向、ジェネリック、命令型など広く対応しており、多目的に使用されています。

  • MacOS(OSX)

    2037questions

    MacOSとは、Appleの開発していたGUI(グラフィカルユーザーインターフェース)を採用したオペレーションシステム(OS)です。Macintoshと共に、市場に出てGUIの普及に大きく貢献しました。

  • 並列処理

    39questions

    複数の計算が同時に実行される手法