Должен ли я объявить двойной массив с номером блока графического процессора во внутреннем или внешнем измерении?

Question

Должен ли я объявить двойной массив с номером блока графического процессора во внутреннем или внешнем измерении?

Например, я должен сделать

int payload[LEN][BLOCKS];

или

int payload[BLOCKS][LEN];

где LEN - очень большое число.

Я планирую, чтобы каждый блок проходил через двойной массив, сохраняя размерность блока постоянной и итерируя по измерению LEN.

Answer 1

Предполагая, что вы собираетесь получать доступ к данным в блочном режиме, вы хотите сделать последнее. Вероятно, это связано с тем, что, когда вы загружаете первый элемент измерения «len», вы уже заплатили за отсутствие в кэше для последующих 7-ми элементов. В первом варианте, вероятно, существует совместное использование строк кэша между блоками графического процессора, но совместное использование является относительно ограниченным и не столь низким.

Действительно, приведенный ниже код сообщает, что для выполнения второго параметра требуется 0,481 секунды, а для первого - 0,979 секунды. Размещение данных с помощью блока по внешнему измерению примерно вдвое быстрее.

#include <cuda_runtime_api.h>
#include <cuda.h>

#include <string>
#include <chrono>
#include <iostream>

#define BLOCKS 80
#define LEN (1 << 20)

void CheckCudaErrorAux (const char *file, unsigned line, const char *statement, cudaError_t err) {
    if (err == cudaSuccess)
        return;
    std::cerr << statement<<" returned " << cudaGetErrorString(err) << "("<<err<< ") at "<<file<<":"<<line << std::endl;
    exit (1);
}
#define CUDA_CHECK_RETURN(value) CheckCudaErrorAux(__FILE__,__LINE__, #value, value)

struct Data1 {
    int payload[LEN][BLOCKS];
};

struct Data2 {
    int payload[BLOCKS][LEN];
};


__global__ void f1(Data1 * data1) {
    int sum = 0;
    for (int i = 0; i < LEN; ++i) {
        sum += data1->payload[i][blockIdx.x];
    }
    printf("block %i has f1 sum %i\n", blockIdx.x, sum);
}

__global__ void f2(Data2 * data2) {
    int sum = 0;
    for (int i = 0; i < LEN; ++i) {
        sum += data2->payload[blockIdx.x][i];
    }
    printf("block %i has f2 sum %i\n", blockIdx.x, sum);
}


int main() {

    Data1 * data1 = (Data1 *) malloc(sizeof(Data1));
    Data2 * data2 = (Data2 *) malloc(sizeof(Data2));;

    for (int i = 0; i < LEN; ++i) {
        for (int b = 0; b < BLOCKS; ++b) {
            data1->payload[i][b] = i * b;
            data2->payload[b][i] = i * b;
        }
    }

    Data1 * data1_on_gpu;
    CUDA_CHECK_RETURN(cudaMalloc(&data1_on_gpu, sizeof(Data1)));
    Data2 * data2_on_gpu;
    cudaMalloc(&data2_on_gpu, sizeof(Data2));
    CUDA_CHECK_RETURN(cudaDeviceSynchronize());
    cudaMemcpy(data1_on_gpu, data1, sizeof(Data1), cudaMemcpyHostToDevice);
    cudaMemcpy(data2_on_gpu, data2, sizeof(Data1), cudaMemcpyHostToDevice);
    CUDA_CHECK_RETURN(cudaDeviceSynchronize());


    std::chrono::time_point<std::chrono::system_clock> t1 = std::chrono::system_clock::now();

    f1<<<80,1>>>(data1_on_gpu);
    CUDA_CHECK_RETURN(cudaDeviceSynchronize());
    std::chrono::time_point<std::chrono::system_clock> t2 = std::chrono::system_clock::now();

    f2<<<80,1>>>(data2_on_gpu);
    CUDA_CHECK_RETURN(cudaDeviceSynchronize());
    std::chrono::time_point<std::chrono::system_clock> t3 = std::chrono::system_clock::now();


    std::chrono::duration<double> duration_1_to_2 = t2 - t1;
    std::chrono::duration<double> duration_2_to_3 = t3 - t2;
    duration_1_to_2.count();

    printf("timer for 1st took %.3lf\n", duration_1_to_2.count());
    printf("timer for 2nd took %.3lf\n", duration_2_to_3.count());

}