1

AMD OpenCL プログラミング ガイド、セクション 6.3 定数メモリの最適化:

グローバル スコープの定数配列。これらの配列は、初期化され、グローバル スコープであり、一定のアドレス空間にあります (OpenCL 仕様のセクション 6.5.3 で指定されています)。配列のサイズが 64 kB 未満の場合、配列はハードウェア定数バッファーに配置されます。それ以外の場合は、グローバル メモリを使用します。この例は、数学関数のルックアップ テーブルです。

この「グローバルスコープの定数配列」を使いたいです。私は純粋なCでそのようなコードを持っています

#define SIZE 101
int *reciprocal_table;

int reciprocal(int number){
  return reciprocal_table[number];
}

void kernel(int *output)
{
  for(int i=0; i < SIZE; i+)
    output[i] = reciprocal(i);
}

OpenCLに移植したい

__kernel void kernel(__global int *output){
  int gid = get_global_id(0);

  output[gid] = reciprocal(gid);
}

int reciprocal(int number){
  return reciprocal_table[number];
}

グローバル変数をどうすればいいですreciprocal_tableか? __globalまたは追加しようとすると__constant、エラーが発生します。

global variable must be declared in addrSpace constant

__constant int *reciprocal_tableからkernelに渡したくありませんreciprocal。どういうわけかグローバル変数を初期化することは可能ですか? コードに書き留めることができることは知っていますが、他の方法はありますか?

PS私はAMD OpenCLを使用しています

UPD上記のコードは単なる例です。多くの機能を備えた、はるかに複雑なコードがあります。そのため、プログラムスコープで配列を作成して、すべての関数で使用したいと考えています。

UPD2サンプル コードを変更し、プログラミング ガイドからの引用を追加

4

3 に答える 3

4 
#define SIZE 2
int constant array[SIZE] = {0, 1};

kernel void
foo (global int* input,
     global int* output)
{
    const uint id = get_global_id (0);
    output[id] = input[id] + array[id];
}

上記をIntelとAMDでコンパイルすることができます。配列の初期化なしでも機能しますが、配列の内容がわからず、定数アドレス空間にあるため、値を割り当てることができませんでした。

標準のセクション 6.5.3 で述べられているように、プログラムのグローバル変数は __constant アドレス空間になければなりません。

更新さて、私は質問を完全に理解しました：

私にとってうまくいったことの1つは、定数空間で配列を定義し、配列をconstant int* array上書きするカーネルパラメーターを渡して上書きすることです。 これにより、GPU デバイスでのみ正しい結果が得られました。AMD CPU デバイスと Intel CPU デバイスは、アレイ アドレスを上書きしませんでした。また、標準に準拠していない可能性もあります。

外観は次のとおりです。

#define SIZE 2
int constant foo[SIZE] = {100, 100};

int
baz (int i)
{
  return foo[i];
}

kernel void
bar (global int* input,
     global int* output,
     constant int* foo)
{
    const uint id = get_global_id (0);
    output[id] = input[id] + baz (id);
}

入力 = {2, 3} および foo = {0, 1} の場合、HD 7850 デバイス (Ubuntu 12.10、Catalyst 9.0.2) で {2, 4} が生成されます。しかし、CPU では、いずれかの OCL 実装 (AMD、Intel) で {102, 103} を取得します。ですから、これが破綻するのは時間の問題なので、私が個人的にどれだけこれをしないかを強調することはできません.

これを実現する別の方法は、実行時に配列の定義 (または事前定義) を使用してホストで .h ファイルを計算し、コンパイラ オプションを介してコンパイル時にカーネルに渡すことです。もちろん、これには異なる LUT ごとに clProgram/clKernel の再コンパイルが必要です。

于 2013-07-02T21:51:32.143 に答える
1

少し前に、自分のプログラムでこの作業を行うのに苦労しました。clEnqueueWriteBuffer などを介して、ホストから定数またはグローバル スコープ配列を初期化する方法が見つかりませんでした。唯一の方法は、.cl ソース ファイルに明示的に記述することです。

したがって、ここでホストから初期化するための私のトリックは、実際にホストからソースをコンパイルしているという事実を利用することです。これは、コンパイルする前に src.cl ファイルを変更できることも意味します。

最初に私の src.cl ファイルは次のようになります。

__constant double lookup[SIZE] = { LOOKUP };    // precomputed table (in constant memory).

double func(int idx) {
  return(lookup[idx])
}

__kernel void ker1(__global double *in, __global double *out)
{
   ... do something ...
   double t = func(i)
   ...
}

ルックアップ テーブルが LOOKUP で初期化されていることに注意してください。

次に、ホスト プログラムで、OpenCL コードをコンパイルする前に、次のようにします。

  • host_values[] でルックアップ テーブルの値を計算する

  • ホストで、次のように実行します。

    char *buf = (char*) malloc( 10000 );
int count = sprintf(buf, "#define LOOKUP ");    // actual source generation !
for (int i=0;i<SIZE;i++) count += sprintf(buf+count, "%g, ",host_values[i]);
count += sprintf(buf+count,"\n");

  • 次に、ソース ファイル src.cl の内容を読み取り、buf+count に配置します。

  • これで、ホストから計算したルックアップ テーブルが明示的に定義されたソース ファイルが作成されました。
  • clCreateProgramWithSource(context, 1, (const char **) &buf, &src_sz, err); のようなものでバッファをコンパイルします。
  • ほら！
于 2013-07-03T17:48:53.967 に答える
-1

「配列」は一種のルックアップテーブルのようです。GPU がそのコピーを使用できるようにするには、clCreateBuffer と clEnqueueWriteBuffer が必要です。

于 2013-07-02T21:59:04.400 に答える