0

問題

プログラムが CUDA カーネルの特定のブランチに到達した回数を数える最良の方法を見つけようとしています。一部のイベントはほとんど発生しないという考えですが、GPU によって処理されたデータは数値最適化ソルバーによって与えられるため、不明確なケースがより一般的になる状況が発生する可能性があります。したがって、複数のシミュレーションでこれらの現象を追跡/監視して、後でグローバルな統計を作成できるようにしたいと考えています。

考えられるアイデア

これを行う最も簡単な方法は、そのような発生を監視する専用の構造を使用することです。次に、監視対象のブランチに入ると、 を使用して関連するカウンターをインクリメントしますatomicAdd。シミュレーションの最後に、カウンターをホストにコピーし、将来の統計処理のために保存します。

私の場合、これらのブランチにあまり入る必要がないため、使用コストはatomicAddそれほど重要ではありませんが、後で一般的なブランチのいくつかを監視したい場合があるので、より良いアプローチは何でしょうか? これは監視のためだけのものなので、オーバーヘッドをあまり重要視したくありません。

ブロックごとに 1 つの監視構造を持ち、最後に合計を実行することもできると思います。これは、グローバル メモリをあまり使用しないためです (unsigned int監視対象のブランチごとに 1 つ)。

コード例

#include <iostream>
#include <time.h>
#include <cuda.h>
#include <stdio.h>

#define CUDA_CHECK_ERROR()  __cuda_check_errors(__FILE__, __LINE__)
#define CUDA_SAFE_CALL(err) __cuda_safe_call(err, __FILE__, __LINE__)

inline void __cuda_check_errors(const char *filename, const int line_number)
{
    cudaError err = cudaDeviceSynchronize();
    if(err != cudaSuccess)
    {
        printf("CUDA error %i at %s:%i: %s\n",
               err, filename, line_number, cudaGetErrorString(err));
        exit(-1);
    }
}

inline void __cuda_safe_call(cudaError err, const char *filename, const int line_number)
{
    if (err != cudaSuccess)
    {
        printf("CUDA error %i at %s:%i: %s\n",
               err, filename, line_number, cudaGetErrorString(err));
        exit(-1);
    }
}

struct Stats
{
    unsigned int even;
};

__global__ void test_kernel(int* A, int* B, Stats* stats)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;

    int res = A[tid] + (int)tid;

    if (res%2 == 0)
        atomicAdd(&(stats->even), 1);

    B[tid] = res;
}

int get_random_int(int min, int max)
{
    return min + (rand() % (int)(max - min + 1));
}

void print_array(int* ar, unsigned int n)
{
    for (unsigned int i = 0; i < n; ++i)
        std::cout << ar[i] << " ";
    std::cout << std::endl;
}

void print_stats(Stats* s)
{
    std::cout << "even: " << s->even << std::endl;
}

int main()
{
    // vector size
    const unsigned int N = 10;

    // device vectors
    int *d_A, *d_B;
    Stats *d_stats;

    // host vectors
    int *h_A, *h_B;
    Stats *h_stats;

    // allocate device memory
    CUDA_SAFE_CALL(cudaMalloc(&d_A, N * sizeof(int)));
    CUDA_SAFE_CALL(cudaMalloc(&d_B, N * sizeof(int)));
    CUDA_SAFE_CALL(cudaMalloc(&d_stats, sizeof(Stats)));

    // allocate host memory
    h_A = new int[N];
    h_B = new int[N];
    h_stats = new Stats;

    // initialize host data
    srand(time(NULL));
    for (unsigned int i = 0; i < N; ++i)
    {
        h_A[i] = get_random_int(0,10);
        h_B[i] = 0;
    }
    memset(h_stats, 0, sizeof(Stats));

    // copy data to the device
    CUDA_SAFE_CALL(cudaMemcpy(d_A, h_A, N * sizeof(int), cudaMemcpyHostToDevice));
    CUDA_SAFE_CALL(cudaMemcpy(d_stats, h_stats, sizeof(Stats), cudaMemcpyHostToDevice));

    // launch kernel
    dim3 grid_size, block_size;
    grid_size.x = N;
    test_kernel<<<grid_size, block_size>>>(d_A, d_B, d_stats);

    // copy result back to host
    CUDA_SAFE_CALL(cudaMemcpy(h_B, d_B, N * sizeof(int), cudaMemcpyDeviceToHost));
    CUDA_SAFE_CALL(cudaMemcpy(h_stats, d_stats, sizeof(Stats), cudaMemcpyDeviceToHost));

    print_array(h_B, N);
    print_stats(h_stats);

    // free device memory
    CUDA_SAFE_CALL(cudaFree(d_A));
    CUDA_SAFE_CALL(cudaFree(d_B));
    CUDA_SAFE_CALL(cudaFree(d_stats));

    // free host memory
    delete [] h_A;
    delete [] h_B;
    delete h_stats;
}

ハードウェア/ソフトウェア情報

私が探しているソリューションは、CC >= 2.0デバイスとCUDA >= 5.0.

4

1 に答える 1