Несвязные магазины

Question

Почему это ядро ​​создает несвязные хранилища

__global__ void reverseArrayBlock(int *d_out, int *d_in)
{
    int inOffset  = blockDim.x * blockIdx.x;
    int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);
    int in  = inOffset + threadIdx.x;
    int out = outOffset + (blockDim.x - 1 - threadIdx.x);
    d_out[out] = d_in[in];
}

а этот нет

__global__ void reverseArrayBlock(int *d_out, int *d_in)
{
    extern __shared__ int s_data[];

    int inOffset  = blockDim.x * blockIdx.x;
    int in  = inOffset + threadIdx.x;

    // Load one element per thread from device memory and store it 
    // *in reversed order* into temporary shared memory
    s_data[blockDim.x - 1 - threadIdx.x] = d_in[in];

    // Block until all threads in the block have written their data to shared mem
    __syncthreads();

    // write the data from shared memory in forward order, 
    // but to the reversed block offset as before

    int outOffset = blockDim.x * (gridDim.x - 1 - blockIdx.x);

    int out = outOffset + threadIdx.x;
    d_out[out] = s_data[threadIdx.x];
}

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

Answer 1

Объединение требует, чтобы адреса следовали шаблону "base + tid" внутри деформации, где tid является коротким для индекса потока. Другими словами, с увеличением tid увеличивается и адрес. Ваш комментарий называет это "предварительным заказом". В первом ядре адреса генерируются так, что с увеличением tid адрес уменьшается, то есть доступ осуществляется в «обратном порядке».

Answer 2

Прежде чем мы начнем, вам нужно понять, что запись в общую память намного дешевле, чем запись в глобальную память.

Имея это в виду, допустим, что мы инвертируем массив 1-> 32

метод первый делает это: во время записи поток 1 читает из местоположения x, поток 2 читает из (x + 1), поток 3 читает из местоположения (x + 2) ... поток 32 читает из местоположения (x + 31).

Вы можете прочитать весь этот кусок памяти в 2 (если выровнены) или 3 (если не выровнены) чтения, потому что операции выполняются в виде фрагментов полусверток (16 потоков).

во время письма поток 1 записывает в местоположение (y + 31), поток 2 записывает в (y + 30), поток 3 записывает в расположение (y + 29) ... поток 32 записывает в местоположение (y).

Хотя они пишут в непрерывный кусок памяти, они находятся в обратном порядке. Если вы не используете какое-то новейшее оборудование (и даже с этим я сомневаюсь), для выполнения операции потребуется 32 записи.

Что касается второго случая, вы выполняете 32 обратных записи в общую память и 32 обратных чтения из общей памяти, что не так дорого.

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

поток 1 записывает в местоположение y, поток 2 записывает в местоположение y + 1 и так далее.

Суть в том, что вы экономите время, затрачиваемое на выполнение 32 (метод 1) - 3 (метод 2) = 29 операций записи в глобальную память.