Как поделиться одним массивом с несколькими блоками и потоками?

Question

Я пытаюсь выучить CUDA и пытаюсь пройти простую программу.Программа просматривает предварительно заполненный массив, заполненный 0,1,2, а затем подсчитывает количество связанных чисел в общем массиве (т. Е. Сколько 00,01,02,10,11,12,20,21,22комбинаций).К сожалению, кажется, что он считает только 1 из каждого вхождения, а затем останавливается.

Предварительно заполненный массив имеет (0,1,0,2,0,0,2,0,1,0) Ожидаемый результат долженbe (1,2,2,2,2,0,0,2,0,0) Фактический объем производства составляет (1,1,1,1,1,0,0,1,0,0)

int a * - это предварительно заполненный массив, int b * - это «общий» массив комбинаций.

В настоящее время глобальный ядро ​​вызывается с одним блоком из 10 потоков.(Позже я хотел бы изменить это на несколько блоков, но сначала я хотел, чтобы потоки работали).

Есть предложения?

Я пытался поделиться массивом, используя shared и использование __syncthreads, но моя проблема, вероятно, требует чего-то другого.

__device__ int GetIndex(int a, int b) {
    if (a == 0 && b == 0) return 0;
    if (a == 0 && b == 1) return 1;
    if (a == 0 && b == 2) return 2;
    if (a == 1 && b == 0) return 3;
    if (a == 1 && b == 1) return 4;
    if (a == 1 && b == 2) return 5;
    if (a == 2 && b == 0) return 6;
    if (a == 2 && b == 1) return 7;
    if (a == 2 && b == 2) return 8;
}

__global__ void CalculateRecurrences(int *a, int *b) {

    __shared__ int s[TOTAL_COMBINATIONS];
    int e = threadIdx.x + blockIdx.x * blockDim.x;

    for (int i = 0; i < 10; i++)
    {
        s[i] = b[i];        
    }
    __syncthreads();
    if (e < 10) {
        int index;
        int next = a[e + 1];
        printf("%i %i", a[e], next);
        index = GetIndex(a[e], next);
        s[index] += 1;
    }

    for (int i = 0; i < 10; i++)
    {
        b[i] = s[i];        
    }
    __syncthreads();
}

Заранее спасибо.Пожалуйста, дайте мне знать, если мне нужно что-то уточнить.

Answer 1

Здесь есть ряд проблем.

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

for (int i = 0; i < 10; i++)
{
    s[i] = b[i];        
}

выполняются всеми 10 вашими потоками.Таким образом, все 10 потоков читают и записывают все 10 элементов вашего входного массива в ваш общий массив.Как расточительно!У вас есть 10 потоков, и у вас есть 10 элементов;вы можете просто указать каждому потоку работать с одним из элементов, заменив указанный выше цикл for на:

if (e < 10)
    s[e] = b[e];

Аналогично, у вас есть 10 потоков, которые все пытаются выполнить следующий блок кода.Вы обращаетесь к памяти способом, который не является потокобезопасным.Самым простым решением было бы использовать atomicAdd вместо + =.

У вас также есть незаконный доступ к памяти здесь;если a определено в диапазоне 0-9, а e определено в диапазоне 0-9, то e+1 превысит границы a:

int next = a[e + 1];  // Undefined behavior!!!

Наконец, как и вышеу вас есть каждый поток, выполняющий ваш последний цикл для копирования элементов из s в b.У вас есть 10 потоков и 10 элементов для работы, поэтому вы должны позволить каждому потоку работать со своим соответствующим индексом:

b[e] = s[e]; 

edit: Соберите все вместе, ваш код может выглядеть примерно так (не проверено):

__global__ void CalculateRecurrences(int *a, int *b) {    
    int e = threadIdx.x + blockIdx.x * blockDim.x;
    if (e >= 10) {
        return;
    }

    __shared__ int s[10];

    // All threads read and assign a value to shared memory
    s[e] = 0;  // You're counting; assign 0s everywhere, we don't need array b for this
    // Wait for all threads to complete initialization of shared array
    __syncthreads();

    // Each thread compares its indexed value to the value in the next index
    if (e < 9) {
        int next = a[e + 1];  
        printf("%i %i", a[e], next);
        int index = GetIndex(a[e], next);
        // Since multiple threads may receive same index, need atomicAdd:
        atomicAdd(&s[index], 1);
    }
    // Each thread may be updating different indices than its own.
    // Thus need to wait for all threads to complete
    __syncthreads();

    // Each thread writes its indexed value to global output array
    b[e] = s[e];
}

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