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

Question

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

Вот мое ядро:

__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-массивах, если вы только начинаете изучать куду.

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

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

1. при использовании 2D массивов, это почти всегда необходимо передать параметр шаг (в некотором роде) в ядре. cudaMallocPitch выделяет массивы с дополнительным заполнением в конце каждой строки, так что следующий ряд начинается с красиво выровненной границы. Обычно это приводит к гранулярности выделения 128 или 256 байтов. Итак, ваша первая строка имеет 3 действительных объекта данных, за которыми следует достаточно пустого места, чтобы заполнить вверх, скажем, 256 байт (равный любой переменной основного тона). Поэтому мы должны изменить вызов ядра и само ядро ​​для учета этого.
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; 
} 

Вы также можете найти this question интересное чтение.

EDIT: отвечая на вопрос в комментариях:

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

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

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

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

Answer 2

На самом деле это может быть сделано путем замены линии

int width_in_bytes = 3 * sizeof(float); 

по:

int width_in_bytes = sizeof(float)*9; 

, потому что это параметр, который говорит cudaMemcpy2D, сколько байт для копирования с СРК в целевой_адрес, в первый код, который вы запрашиваете, чтобы скопировать 3 числа с плавающей точкой, но массив, который вы хотите скопировать, имеет длину 9, поэтому требуемая ширина - это 9 чисел с плавающей запятой.

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

threadIdx.y==0