問題タブ [reduction]

私は OpenCl の初心者です。

double の 1 次元配列に対してリダクション (合計演算子) を操作する必要があります。

私はネットを徘徊してきましたが、見つけた例はかなり混乱しています。誰でも読みやすい (そしておそらく効率的な) チュートリアルの実装を投稿できますか?

追加情報: - 1 つの GPU デバイスにアクセスできます。- カーネル コードに C を使用しています

次のコードは、amd Web サイトからのものです。

合計削減として機能するように調整しました。

そして、ワークグループを 1 つしか使用しない場合 (つまり、 に与える) は魅力的に機能NULLしlocal_work_sizeますclEnqueueNDRangeKernel()が、ワークグループの次元を変更しようとすると、制御不能になります。(私は OpenCl の初心者です)

私がすることは次のとおりです

やり方が悪いのでしょうか？？

#define GLOBAL_DIM 60000さらに、設定すると（これが必要になります）、ローカルメモリが不足することに気付きました。複数のワーク グループを使用すると、「より多くの」ローカル メモリを取得できますか?

並列削減を実行したいのですが、共有メモリ内のデータを使用してカーネル内で実行します。これはスラストライブラリで可能ですか？何かのようなもの

しかし、これはカーネル内では機能しません。出来ますか？ありがとうございました。

このコードで配列を合計しようとしていますが、スタックしています。このような基本的な操作に多くの時間を費やし、機能させることができないため、おそらく「CUDA for dummies チュートリアル」が必要です。

以下は、私が理解できないこと、または確信が持てないことのリストです。

1. 何個のブロック (dimGrid) を使用する必要がありますか? N/dimBlock.x/2カーネルの最初に、グローバルメモリの2つの「ブロック」からデータがロードされ、共有メモリに追加されるため、（N =入力配列の長さ）である必要があると思います

2. 元のコードにはblockSize. blockDim.xこれらの変数の違いがわからないので、に置き換えました。しかし、blockSize=blockDim.xの場合、gridSize = blockDim.x*2*gridDim.x私には意味がありませんgridSize。N よりも大きくなります。1D 配列のコンテキストでの *Dim.x と *Size の違いは何ですか?

3. 主なロジック - カーネルでは、各ブロックは 2*dimBlock(ブロック内のスレッド) の数値を合計します。N = 262144 および dimBlock = 128 の場合、カーネルは部分和の 1024 配列を返します。次に、カーネルを再度実行すると、結果 = 4 つの部分合計が得られます。最後に、最後の実行では、配列が単一のブロックで処理されるため、単一の合計が返されます。

4. バイナリ配列を合計します。最初の実行ではuchar4、入力データに使用できます。2 回目と 3 回目の実行では、 を使用しますint。

何が欠けているのか教えてください

ありがとう

以下のコードでは、エラーが発生します：unrecognized #pragma: #pragma omp reduction (+: sum)。関数自体はすでに並列化されているため、関数内のforループはparallel-for-loopではないことに注意してください。問題はどこにあるのでしょうか？

メインCPPファイル：

関数が定義されている別のcppファイル

システムに UUID を保存するために QR コード バーコードを使用していますが、生成されたバーコードが自分のものであり、他人のものではないことを確認する必要があります。また、QR コードが下位バージョンの範囲に留まり、スキャンしやすいままになるように、エンコードされたデータを短く保つ必要もあります。

私のアプローチは、UUID の raw 値の数値 (128 ビット値) と 16 ビットのチェックサムを取得し、そのデータを Base64 でエンコードしてから QR コードに変換することです。これまでのところ、これは完全に機能します。

チェックサムを生成するには、UUID の文字列バージョンを取得し、それを長い秘密の文字列と結合し、奇数バイトを XOR して SHA-1 ハッシュを生成します。しかし、このハッシュは長すぎるので、すべての古いバイトを XOR してチェックサムの半分を生成し、同様に偶数バイトを使用して残りの半分を生成します。

私が心配しているのは、SHA-1 システムを不必要に XOR ダウンして侵害したことです。結果内のどこかから、操作されていない 2 バイトだけを取得する方がよいでしょうか? 16 ビットのチェックサムが 160 ビットのチェックサムほど安全ではないことは承知していますが、それはバーコードを使いやすくするために支払わなければならない代償です。私が本当に見つけたくないのは、UUIDが平文で送信されるため、簡単にクラックできるチェックサムを提供したことです。

チェックサムを生成するより良い方法があれば、それも質問に対する適切な答えになります。いつもお時間をいただき、またはこれを読んでいただき、ありがとうございます。回答を投稿していただければ、さらに感謝いたします。

次の情報を含むデータセットがあります。

• ワークロード名
• 使用した構成
• 測定されたパフォーマンス

ここに、私の問題を説明するおもちゃのデータ セットがあります (パフォーマンス データはまったく意味がありません。例を簡単に理解できるように、さまざまな整数を選択しただけです。実際には、そのデータはパフォーマンス測定から得られた浮動小数点値になります)。

次を使用して生成できます。

さまざまな構成の高調波スピードアップを計算しようとしています。そのためには、基本構成が必要です (この例では cfg = 1)。次に、ハーモニック スピードアップは次のように計算されます。

たとえば、構成 2 の場合は次のようになります。

すべてのワークロード ペアと構成の高調波スピードアップを計算したいと考えています。サンプル データ セットを使用すると、結果は次のようになります。

ループを使用しない解決策を見つけるために苦労してaggregateいddplyますが、実用的な解決策を思いつくことができませんでした。したがって、私が直面している基本的な問題は次のとおりです。

• ワークロードと構成の間の関係を処理する方法。特定のワークロード ペア (AB) と特定の構成の結果は、一緒に処理する必要があります (ハーモニック スピードアップ式の分母の最初の 2 つのパフォーマンス測定値はワークロード A から取得され、残りの 2 つはワークロード B から取得されます)。
• ワークロードのペアと構成ごとに、構成ベース (例では cfg 1) の値を使用してパフォーマンス値を「正規化」する必要があります。

aggregateorなどのR関数でそれを表現する方法がよくわかりませんddply（可能であれば、まったく）。

これを解決する方法を知っている人はいますか？

編集: 1..8 as を使用perfすると混乱が生じるのではないかと心配していました。簡単にするためにそうしましたが、値はそれらの値である必要はありません (たとえば、次のように初期化することを想像してください: dframe$perf <- runif(8))。ジェームズとザックの両方の回答は、私の質問のその部分が間違っていることを理解していたので、質問でこれを明確にする方がよいと思いました. とにかく、構成1のパフォーマンスが（1、2）ではない場合に対処するために、両方の答えを一般化しました

ワープのすべてのスレッドは定義上同期しているので、ワープベースの並列削減について考えました。

したがって、同期を必要とせずに、入力データを64分の1に減らすことができる（各スレッドは2つの要素を減らす）という考え方でした。

Mark Harrisによる元の実装と同じように、削減はブロックレベルで適用され、データは共有メモリにあります。 http://gpgpu.org/static/sc2007/SC07_CUDA_5_Optimization_Harris.pdf

彼のバージョンとワープベースのバージョンをテストするためにカーネルを作成しました。
カーネル自体は、BLOCK_SIZE要素を共有メモリに完全に同じように格納し、その結果を出力配列の一意のブロックインデックスに出力します。

アルゴリズム自体は正常に機能します。「カウント」をテストするために、自分の完全な配列でテストされました。

実装の関数本体：

1.彼のバージョンの実装：

リソース：

4つの同期スレッドが使用されました12if
ステートメントが使用されまし
た11読み取り+追加+書き込み操作
1最終書き込み操作
5レジスタの使用

パフォーマンス：

5回のテスト実行の平均：〜19.54ミリ秒

2.ワープベースのアプローチ:(上記と同じ機能本体）

リソース：

1つのsyncthreadが使用されました7if
ステートメント
10読み取り追加書き込み操作
1最終書き込み操作
5レジスタの使用

5ビットシフト
1追加
1サブ

パフォーマンス：

5回のテスト実行の平均：〜20.82ミリ秒

256mbのfloat値を持つGeforce8800GT512mbで両方のカーネルを複数回テストします。そして、ブロックあたり256スレッドでカーネルを実行します（100％の占有率）。

ワープベースのバージョンは、約1.28ミリ秒遅くなります。

将来のカードでより大きなブロックサイズが許可される場合、最大値が4096で64に減少し、最終的なワープによって1に減少するため、ワープベースのアプローチではそれ以上の同期ステートメントは必要ありません。

なぜそれは速くないのですか？またはアイデアの欠陥はどこにありますか、カーネル？

リソースの使用法から、ワープアプローチを先に進める必要がありますか？

編集1：スレッドの半分だけがアクティブであり、範囲外の読み取りが発生しないようにカーネルを修正し、新しいパフォーマンスデータを追加しました

剰余累乗の一般的な方程式は、(a + b) MOD n = ((a MOD n) + (b MOD n)) MOD n です。a と b が非常に大きい場合、これはすばらしいことです。ただし、この累乗を非常に大きな n (2^31 -1) で行うように求められましたが、a と b は問題ありません。

nを減らす方法が必要です。

以下は、削減するのが難しいと感じているラムダ式です。つまり、この問題の対処方法を理解できません。

(λmn(λsz.ms(nsz)))(λsz.sz)(λsz.sz)

私はこれで迷っています。

誰かが私を正しい方向に導くことができれば、それは大歓迎です