CUDA Vector Reduction для обработки векторов длиной менее 512? - PullRequest
Новая
       78

CUDA Vector Reduction для обработки векторов длиной менее 512?

0 голосов
/ 19 февраля 2019

Я работаю над учебным пособием по параллельному алгоритму vector_reduction от NVIDIA, чтобы реализовать алгоритм с использованием CUDA C ++ API.Я реализовал алгоритм, но он работает только для векторной длины, которая установлена ​​на 512. Я не могу понять, как заставить его работать для векторов, меньших 512?Я хочу, чтобы он работал для произвольных размеров, то есть 324, 123, 23. 

#include <stdio.h>

#define NUM_ELEMENTS 512

__global__ void reduction(float *g_data, int n)
{
    __shared__ float partialSum[NUM_ELEMENTS];

    int tx = threadIdx.x;
    int i = tx + blockIdx.x * blockDim.x;

    if (i < n) {
        partialSum[tx] = g_data[i];
    }

    int stride;
    for (stride = blockDim.x/2; stride > 0;  stride >>= 1) {
        __syncthreads();
        if (tx < stride) {
           partialSum[tx] += partialSum[tx + stride];
        }
    }

    if (tx == 0) {
        g_data[blockIdx.x] = partialSum[tx];
    }
}

float computeOnDevice(float* h_data, int num_elements)
{
    float* d_data = NULL;
    float result;

    // Memory allocation on device side
    cudaMalloc((void**)&d_data, sizeof(float)*num_elements);

    // Copy from host memory to device memory
    cudaMemcpy(d_data, h_data, num_elements * sizeof(float), cudaMemcpyHostToDevice );

    dim3 blockSize, gridSize;

    // Number of threads in each thread block
    blockSize = dim3(num_elements, 1, 1);

    // Number of thread blocks in grid
    gridSize = dim3(1, 1, 1);

    // Invoke the kernel
    reduction<<<gridSize, blockSize>>>(d_data, num_elements);

    // Copy from device memory back to host memory
    cudaMemcpy(&result, d_data, sizeof(float), cudaMemcpyDeviceToHost);
    cudaFree(d_data);
    cudaDeviceReset();
    return result;
}

int main() {

    float *data = new float[NUM_ELEMENTS];
    for (int i = 0; i < NUM_ELEMENTS; i++) data[i] = 1;
    float r = computeOnDevice(data, NUM_ELEMENTS);
    printf(" result = %f\n" , r);
}

1 Ответ

0 голосов
/ 19 февраля 2019

Ваш код на 100% правильный.Проблема в том, что ваши битовые сдвиги не учитывают последнюю часть вашего массива.Вы можете легко исправить это, искусственно расширив массив до следующей степени 2. Таким образом, весь ваш массив будет уменьшен, а лишние «элементы» (они на самом деле не существуют) просто игнорируются. 

#include <math.h>

__global__ void reduction(float *g_data, int n){
    // figure out exponent of next larger power of 2
    int exponent = ceilf(log2f(n));
    // calculate next larger power of 2
    int size = (int)powf(2, exponent);
    __shared__ float partialSum[NUM_ELEMENTS];

    int tx = threadIdx.x;
    int i = tx + blockIdx.x * blockDim.x;

    if (i < n){
        partialSum[tx] = g_data[i];
    }

    for (int stride = size / 2; stride > 0; stride >>= 1){
        __syncthreads();

        if (tx < stride) {
            // all threads that run out of bounds do nothing
            // equivalent to adding 0
            if((tx + stride) < n)
                partialSum[tx] += partialSum[tx + stride];
        }
    }

    if (tx == 0){
        g_data[blockIdx.x] = partialSum[tx];
    }
}

Редактировать

Что касается вашего комментария, этот метод сокращения никогда не будет работать для массива, который сокращается в несколько блоков.Таким образом, для вычислительных возможностей 1.0-1.3 самый большой массив, который вы можете уменьшить, составляет 512 элементов, для вычислительных возможностей> 1.3 вы можете сделать до 1024 элементов, это максимальное количество потоков в блоке.

Это потому, что __shared__ память распределена между потоками, а не блоками .Таким образом, чтобы уменьшить массив, разбросанный по нескольким блокам, вам нужно разделить массив так, чтобы каждый блок уменьшал кусок, а затем использовать память __global__ для уменьшения значений из всех блоков.Однако __global__ память примерно в 10-20 раз медленнее, чем (встроенная) __shared__ память, поэтому, как только вы начнете использовать много блоков, это станет очень неэффективным.

Альтернатива будетчтобы каждый поток обрабатывал несколько индексов, однако, в конечном итоге ваш массив partialSum больше не помещается в разделяемую память и в любом случае переполняется в глобальную память.Этот подход также означает, что вы никогда не сможете использовать более 512 (или 1024) потоков, что противоречит цели использования CUDA, которая зависит от запуска очень большого количества потоков, чтобы скрыть задержку и сделать дорогостоящую передачу памяти от хоста к устройству стоящей..

...