Насколько быстро выполняется запрос 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;
}

Без оптимизации nv cc я получил следующие результаты:

1 => 45.5977 ms
2 => 45.4554 ms

, что указывает на та же производительность. Опять же, я не уверен, является ли вывод здесь обобщенным, и подозревает, что это зависит от вашего ядра и nv cc.