1

私は次のことをしたいと思います: サイズ N (N はかなり大きい) と数値 x の並べ替えられた数値ベクトルがあるとしましょう。このベクトル内の番号 x の正しい位置の検索を並列化したいと考えています。例えば：

myVector = [ 1, 2, 3, .... , 10000] および x = 3.2,

次に、3 を返す必要があります。正しい位置を見つけた最初のスレッドは、他のスレッドのジョブを中断する必要があります。t= min(t_1, t_2,......, t_number of threads) 正しい位置を探すためにマルチスレッドを使用する方が速くなると思いますか? スレッド間の通信はどうですか？値がスレッドによって赤色になり、検索に一致しないため、他のスレッドは検索中にこの値をスキップする必要があります (変更する必要があるブール値である可能性があります..

このアルゴリズムに関して何かアドバイスはありますか?

4

2 に答える 2

0

しばらく前に、同様のことを行う次のコードを書きました。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <stdlib.h>

__global__ void fast_finder(unsigned int *g_found, float x, float *y)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int pos = (unsigned int)(x == y[i]);
    g_found[i * (1 - pos)] = i * pos;
}

int main(int argc, char *argv[])
{
    int N = 65536;
    unsigned int h_found, *d_found;
    float *h_y = (float *)malloc(N * sizeof(float)), *d_y, x = 5.0f;
    int nThreads = 1024, nBloks = N / nThreads;

    for (int i = 0; i < N; ++i) h_y[i] = (float)(N - i - 1);

    if (x != h_y[0]) {
        cudaSetDevice(0);
        cudaMalloc((void **)&d_found, N * sizeof(unsigned int));
        cudaMalloc((void **)&d_y, N * sizeof(float));
        cudaMemcpy(d_y, h_y, N * sizeof(float), cudaMemcpyHostToDevice);

        fast_finder<<<nBloks, nThreads>>>(d_found, x, d_y);
        cudaThreadSynchronize();

        cudaMemcpy(&h_found, d_found, sizeof(unsigned int), cudaMemcpyDeviceToHost);
        if (h_found) printf("%g found on %d. position!\n", x, h_found);
        else printf("%g not found!\n", x);

        cudaFree(d_y);
        cudaFree(d_found);

    } else printf("%g found on the first position!\n", x);

    free(h_y);

    getchar();
    return EXIT_SUCCESS;
}

ここで、各スレッドは、グローバル スレッド インデックスによって提供される値yが と等しいかどうかをチェックしxます。true の場合、スレッドはそのインデックスをg_found配列の最初の位置に書き込みます。それ以外の場合はg_found、インデックスによって提供される位置に 0 を書き込みます。y長さ 16 の場合、出力の 11 番目の位置に値 5 を含むとy、次のようになります。

g_found = { 10, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0 }

この場合y、並べ替える必要はありませんが、一意の値のみを含める必要があります。xこのコードは、次のように、提供されたものが挿入される 1 つの検出結果 (デバイス部分) インデックスに簡単に変更できます。

__global__ void fast_finder(unsigned int *g_found, float x, float *y)
{
    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
    unsigned int pos = (unsigned int)(x >= y[i] || x <= y[i+1]);
    g_found[i * (1 - pos)] = (i + 1) * pos;
}

このバージョンの出力は、私のものと似ています。g_found位置 0 が 0 の場合、 の値は配列xに存在しませんy。isの最初の要素yが等しいかどうかはx、カーネルが呼び出される前に、ホスト コードによってチェックされます。この部分も変更して条件を適用しても問題ありません。

ご覧のとおり、このようなソリューションでは、すべてのスレッドが連携して動作し、実行を終了する必要はありませんx。良いことは、パケット検索を適用することです。つまり、1 つのスレッドを の小さなサブセットでシークするように割り当てることで、はるかに大きくすることがyできます。y

于 2013-03-07T10:26:46.900 に答える
0

スレッドとブロックの間で通信する必要はありません。現在のインデックスの値が予想よりも大きいかどうかを確認できます。なら返す。ほとんどのスレッドは、このチェックに耐えられません。

これで、値が期待値より小さいインデックスを持つスレッドのみができました。次の値がクエリ以上かどうかを確認し、適切なインデックスを返します。

これは午前 5 時に書いている未テストのカーネルです。

template<typename ty>
__global___ static void search(int *out, ty *list, ty val, int n)
{
    int start = threadIdx.x + blockIdx.x * blockDim.x;
    for (int idx = start; idx < n; idx += gridDim.x * blockDim.x) {
        if (list[idx] >= val) return;
        ty next = list[idx + 1];
        if (idx == n-1 || next >= val) {
            *out = next == val ? (idx + 1) : idx;
            return;
        }
     }
}

そうは言っても、あなたは本当にこれをしたくありません。CPU の使用中に O(log n) の最悪のケースのパフォーマンスを得ることができます。つまり、10 億個の要素を 32 ステップで検索できるということです。GPU に既にデータがあり、メモリ コピーを避けたい場合を除き、これは CPU の方がはるかに優れています。

于 2013-03-07T10:25:05.127 に答える