0

私のカーネルでは、条件が満たされた場合、出力バッファーの項目を更新します

if (condition(input[i])) //?
    output[i] = 1;

そうしないと、出力が同じままになり、値が0になる可能性があります。

入力によっては、更新の密度はまったく予測できません。さらに、どの出力場所が更新されるかも不明です。（私はそれらを強制するかもしれませんが、場合によっては）

私の質問は、すべての項目を書き込むのが良いのか、合体を達成するのが良いのか、それとも選択的な書き込みを行うのが良いのかということです。

output[i] = condition(input[i]); //?

あなたの発言について話し合っていただけませんか。

4

2 に答える 2

1

合体は、ワープ内の一部のスレッドがロードまたはストアに参加していなくても、参加しているすべてのスレッドが合体の要件を満たしている限り、達成されます。したがって、条件付き書き込みはメモリスループットに影響を与えないはずです。

ただし、条件付き書き込みを行うには、ブランチが含まれるため、追加の命令が含まれる場合があります（これは、たとえば、Eugeneが回答で測定したパフォーマンスの違いを説明する可能性があります）。

于 2012-08-23T01:36:27.370 に答える
1

条件付きセット（オプション1）を実行する私のセットアップカーネルでは、1.727usおよびオプション21.399usで実行されます。これは私のコードです（setConditionalの方が速いです）：

__global__ void conditionalSet(unsigned int* array) {
    if ((threadIdx.x & 3) == 0) {
        array[threadIdx.x] = 1;
    }
}

__global__ void setConditional(unsigned int* array) {
    array[threadIdx.x] = (threadIdx.x & 3) == 0 ? 1 : 0;
}
于 2012-08-22T22:32:18.177 に答える