查询threadIdx，blockIdx，blockDim的速度有多快？

Question

这是一个非常简单的问题。 threadIdx，blockIdx和blockDim的读取速度变量？例如，如果我在内核中需要几次，是否应该先将它们写入本地寄存器，还是直接访问它们无关紧要？

所以，本质上，我问的是threadIdx，blockIdx和blockDim在以下插图的哪个内存中起作用：

Answer 1

正如@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。