0

问题

我试图找到最好的方法来计算我的程序在我的 CUDA 内核的某些特定分支中结束的次数。这个想法是某些事件几乎不应该发生，但由于 GPU 处理的数据是由数值优化求解器给出的，因此在某些情况下，定义不明确的情况可能会变得更加普遍。因此，我希望能够在多个模拟中跟踪/监控这些现象，以便稍后进行一些全局统计。

可能的想法

最直接的方法可能是使用专门用于监控此类事件的结构。然后，当进入一个受监控的分支时，我们使用 增加相关的计数器atomicAdd。在模拟结束时，我们将计数器复制回主机并存储它们以供将来进行一些统计处理。

就我而言，使用成本atomicAdd不应该那么重要，因为我不应该进入那些分支那么多，但是，我可能想稍后监控一些常见的分支，那么什么是更好的方法呢？由于这只是为了监控，我不希望开销太重要。

我想我也可以每个块有一个监控结构并在最后做一个求和，因为它不应该使用太多的全局内存（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

0

atomicAdd 是一种可能性，我可能会走那条路。如果您不使用 atomicAdd 函数调用的结果，编译器将发出一个归约操作，例如RED.E.ADD. 只要没有发生太多冲突，减少就会非常快（我有时实际上会使用它，即使我不需要操作是原子的，因为它比从全局内存中加载值、执行算术运算并保存回全局内存）。
您拥有的第二个选项是使用分析器计数器并使用分析器分析结果。有关详细信息，请参阅Profiler 计数器功能。

于 2013-06-17T09:32:38.083 回答