时间戳

首页 > 解决方案 > 查询 threadIdx、blockIdx、blockDim 有多快？

cuda - 查询 threadIdx、blockIdx、blockDim 有多快？

问题描述

这是一个相当直接的问题。threadIdx读取,blockIdx和blockDim 变量的速度有多快？例如，如果我在内核中需要它们几次，我应该先将它们写入本地寄存器，还是直接访问它们没有关系？

所以，从本质上讲，我是在问下图的哪个记忆中存在threadIdx并blockIdx存在blockDim：

在此处输入图像描述

标签： cuda

解决方案

正如@talonmies 提到的，编译器如何处理变量是不可预测的。尽管如此，我还是在我的系统（Windows 10、CUDA 10.2、Tesla k40）上开发了一个简单的测试来分析编译器在我的情况下关于您的问题的行为。让我们#define nTPB 1024作为每个块的线程数。kernel_1存储threadId.x在t位于寄存器中并t多次读取，而在 中kernel_2，threadId.x每次都直接使用。

// kernel_1 stores threadId.x in t
__global__ void kernel_1(const int N, const int offset, const unsigned *v, unsigned *o) 
{
    unsigned n_ept = (unsigned)(ceil)((double)N / nTPB); // No of elements per thread
    unsigned t = threadIdx.x;
    unsigned t_min = t * n_ept;
    unsigned t_max = (t+1) * n_ept;
    for (unsigned i = t_min; i < t_max; i++) {
        if( i < N)
            o[i] = v[i + offset] + t + t / 2;
    }
}

// kernel_2 does not stores threadId.x
__global__ void kernel_2(const int N, const int offset, const unsigned *v, unsigned *o)
{
    unsigned n_ept = (unsigned)(ceil)((double)N / nTPB); // No of elements per thread
    unsigned t_min = threadIdx.x * n_ept;
    unsigned t_max = (threadIdx.x + 1) * n_ept;
    for (unsigned i = t_min; i < t_max; i++) {
        if (i < N)
            o[i] = v[i + offset] + threadIdx.x + threadIdx.x / 2;
    }
}

以以下方式调用每个函数 100 次，我测量了每个内核的性能：

int main()
{

    int N = 100000;
    int offset = 4;

    std::chrono::high_resolution_clock::time_point cpu_startTime;

    unsigned *h_v = new unsigned[N+offset];
    for (int i = 0; i < N+offset; i++)
        h_v[i] = 10;
    unsigned *d_v;
    unsigned *d_o;
    CHECK_CUDA(cudaMalloc((void **)&d_v, (N+offset) * sizeof(unsigned)));
    CHECK_CUDA(cudaMemcpy(d_v, h_v, (N+offset) * sizeof(unsigned), cudaMemcpyHostToDevice));
    CHECK_CUDA(cudaMalloc((void **)&d_o, N * sizeof(unsigned)));

    dim3 threads(nTPB);

    cpu_startTime = std::chrono::high_resolution_clock::now();
    for(int i = 0; i < 100; i++)
        kernel_1 <<<1, threads >>> (N, offset, d_v, d_o);
    CHECK_CUDA(cudaDeviceSynchronize());
    std::chrono::duration<double> elapsed_data_1 = std::chrono::high_resolution_clock::now() - cpu_startTime;

    cpu_startTime = std::chrono::high_resolution_clock::now();
    for (int i = 0; i < 100; i++)
        kernel_2 <<<1, threads >>> (N, offset, d_v, d_o);
    CHECK_CUDA(cudaDeviceSynchronize());
    std::chrono::duration<double> elapsed_data_2 = std::chrono::high_resolution_clock::now() - cpu_startTime;

    double elapsed_1 = 1000 * elapsed_data_1.count(); // elapsed time in ms
    double elapsed_2 = 1000 * elapsed_data_2.count();

    printf("Elapsed time:\n1 => %g ms\n2 => %g ms\n", elapsed_1, elapsed_2);

    return 0;
}

在没有 nvcc 优化的情况下，我得到了以下结果：

1 => 45.5977 ms
2 => 45.4554 ms

这表示相同的性能。同样，我不确定这里的结论是否可以概括并且怀疑它取决于您的内核和 nvcc。

推荐阅读