CUDA: двумерная индексация массива дает неожиданные результаты

Question

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

Вот мое ядро:

__global__ void doStuff(float* data, float* result)
{
    if (threadIdx.x < 9) // take the first 9 threads
    {
        int index = threadIdx.x;
        result[index] = (float) index;
    }
}

И вот соответствующие части моего main:

#include <stdlib.h>
#include <stdio.h>

int main(void)
{
    /*
        Setup
    */
    float simple[] = {-1.0, -2.0, -3.0, -4.0, -5.0, -6.0, -7.0, -8.0, -9.0};

    float* data_array;
    float* result_array;

    size_t data_array_pitch, result_array_pitch;
    int width_in_bytes = 3 * sizeof(float);
    int height = 3;

    /*
        Initialize GPU arrays
    */
    cudaMallocPitch(&data_array, &data_array_pitch, width_in_bytes, height);
    cudaMallocPitch(&result_array, &result_array_pitch, width_in_bytes, height);

    /*
        Copy data to GPU
    */
    cudaMemcpy2D(data_array, data_array_pitch, simple, width_in_bytes, width_in_bytes, height, cudaMemcpyHostToDevice);

    dim3 threads_per_block(16, 16);
    dim3 num_blocks(1,1);

    /*
        Do stuff
    */
    doStuff<<<num_blocks, threads_per_blocks>>>(data_array, result_array);

    /*
        Get the results
    */
    cudaMemcpy2D(simple, width_in_bytes, result_array, result_array_pitch, width_in_bytes, height, cudaMemcpyDeviceToHost);

    for (int i = 1; i <= 9; ++i)
    {
        printf("%f ", simple[i-1]);
        if(!(i%3))
            printf("\n");
    }

    return 0;
}

Когда я запускаю это, я получаю 0.000000 1.000000 2.00000 для первого ряда и мусор для двух других.

Answer 1

Я не уверен, что сосредоточился бы на 2D-массивах, если вы только начинаете изучать cuda.

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

В любом случае, с вашим кодом есть несколько проблем:

1. при использовании 2D-массивов почти всегда необходимо передавать параметр pitch (некоторым образом) в ядро. cudaMallocPitchвыделяет массивы с дополнительным заполнением в конце каждой строки, так что следующая строка начинается с красиво выровненной границы. Это обычно приводит к гранулярности выделения 128 или 256 байтов. Таким образом, ваша первая строка имеет 3 допустимых объекта данных, за которыми следует достаточно пустого пространства, чтобы заполнить его, скажем, 256 байтов (что равно вашей переменной pitch). Таким образом, мы должны изменить вызов ядра и самого ядра, чтобы учесть это.
2. Ваше ядро ​​по своей сути является 1D ядром (оно не понимает и не использует threadIdx.y, например). Поэтому нет смысла запускать 2D-сетку. Хотя в этом случае это не повредит, это создает избыточность, которая может сбивать с толку и создавать проблемы в других кодах.

Вот обновленный код, показывающий некоторые изменения, которые дадут вам ожидаемые результаты, основанные на вышеупомянутых комментариях:

#include <stdio.h>


__global__ void doStuff(float* data, float* result, size_t dpitch, size_t rpitch, int width)
{
    if (threadIdx.x < 9) // take the first 9 threads
    {
        int index = threadIdx.x;
        result[((index/width)*(rpitch/sizeof(float)))+ (index%width)] = (float) index;
    }
}

int main(void)
{
    /*
        Setup
    */
    float simple[] = {-1.0, -2.0, -3.0, -4.0, -5.0, -6.0, -7.0, -8.0, -9.0};

    float* data_array;
    float* result_array;

    size_t data_array_pitch, result_array_pitch;
    int height = 3;
    int width = 3;
    int width_in_bytes = width * sizeof(float);

    /*
        Initialize GPU arrays
    */
    cudaMallocPitch(&data_array, &data_array_pitch, width_in_bytes, height);
    cudaMallocPitch(&result_array, &result_array_pitch, width_in_bytes, height);

    /*
        Copy data to GPU
    */
    cudaMemcpy2D(data_array, data_array_pitch, simple, width_in_bytes, width_in_bytes, height, cudaMemcpyHostToDevice);

    dim3 threads_per_block(16);
    dim3 num_blocks(1,1);

    /*
        Do stuff
    */
    doStuff<<<num_blocks, threads_per_block>>>(data_array, result_array, data_array_pitch, result_array_pitch, width);

    /*
        Get the results
    */
    cudaMemcpy2D(simple, width_in_bytes, result_array, result_array_pitch, width_in_bytes, height, cudaMemcpyDeviceToHost);

    for (int i = 1; i <= 9; ++i)
    {
        printf("%f ", simple[i-1]);
        if(!(i%3))
            printf("\n");
    }
    return 0;
}

Вы также можете найти этот вопрос интересным чтением.

РЕДАКТИРОВАТЬ: отвечая на вопрос в комментариях:

result[((index/width)*(rpitch/sizeof(float)))+ (index%width)] = (float) index;
              1               2                      3

Чтобы вычислить правильный индекс элемента в переданном массиве, мы должны:

1. Вычислить (виртуальный) индекс строки из индекса потока. Мы делаем это, беря целочисленное деление индекса потока по ширине каждой (не переданной) строки (в элементах, а не в байтах).
2. Умножьте индекс строки на ширину каждой переданной строки. Ширина каждой переданной строки задается переданным параметром, который находится в байтах. Чтобы преобразовать этот параметр байтового байта в параметр основного элемента, мы делим его на размер каждого элемента. Затем, умножив количество на индекс строки, вычисленный на шаге 1, мы теперь проиндексировали правильную строку.
3. Вычислить (виртуальный) индекс столбца из индекса потока, взяв остаток (деление по модулю) индекса потока, деленный на ширину (в элементах). Когда у нас есть индекс столбца (в элементах), мы добавляем его к индексу начала правильной строки, вычисленному на шаге 2, чтобы определить элемент, за который будет отвечать этот поток.

Выше приведено достаточно усилий для относительно простой операции, что является одним из примеров того, почему я предлагаю сначала сосредоточиться на базовых концепциях cuda, а не на тональных массивах. Например, я хотел бы понять, как обрабатывать блоки с 1 и 2-мя нитями, а также с 1 и 2-мерными сетками, прежде чем приступить к работе с массивами с тональностью Смежные массивы являются полезным средством повышения производительности для доступа к 2D-массивам (или 3D-массивам) в некоторых случаях, но они ни в коем случае не являются необходимыми для обработки многомерных массивов в CUDA.

Answer 2

На самом деле это также можно сделать, заменив линию

int width_in_bytes = 3 * sizeof(float);

от:

int width_in_bytes = sizeof(float)*9;

поскольку это параметр, который сообщает cudaMemcpy2D, сколько байт нужно скопировать из src в dst, в первом коде вы просите скопировать 3 числа с плавающей запятой, но массив, который вы хотите скопировать, имеет длину 9, поэтому требуемая ширина равна размеру 9 чисел с плавающей запятой.

Хотя это решение работает, в вашем коде все еще есть некоторые недостатки; например, если вы действительно хотите, чтобы первые 9 потоков блока что-то делали, в 'if' вы должны добавить следующее условие с помощью и (&&)

threadIdx.y==0