VPF::Проблема с CUDA, разные результаты Debug и Release

Bitman

Дата 10.7.2015, 13:03 (ссылка)

(нет голосов)

Загрузка ...

Новичок

Профиль
Группа: Участник
Сообщений: 6
Регистрация: 7.6.2005

Репутация: нет
Всего: нет

Проблема с вроде-бы элементарным кодом на CUDA. Высчитывает среднее арифметическое по массиву.
каждый элемент переводится во float и делится на кол-во элементов, это значение помещается в shared memory.
Затем делается scan, все данные в shared memory складываются, результат в последней ячейке. Потом последняя задача в блоке этот результат
добавляет через atomicAdd к переменной в глобальной памяти.
Всё это замечательно работает при компиляции в Debug режиме. Используется CUDA 7.0 и VisualStudio 10. Как только переключаю на Release,
начинает считать некорректно, а точнее среднее значение получается капельку меньше, чем то, что должно быть, например при среднем 1500 на 0.5-0.7 меньше.
Причём эта капелька намного разная от запуска к запуску. Если в Release в настройках CUDA C/C++ , Device включить параметр Generate GPU debug information = Yes (-G),
то Release код начинает работать корректно, но раза в 2 примерно медленней.
Честно говоря, не знаю уже куда копать :(
Вот этот код:
P.S. Могу выслать проект целиком.

Код


const UINT AVG2_THREADS_IN_BLOCK=256;
// кол-во задач в блоке должно быть степенью двойки.

__global__
void average_kernel_WORDS2(const WORD* const d_in_data, UINT data_size, float* const d_out_value) {
    // d_in_data - входные данные в формате 16-ти разрядных беззнаковых целых
    // data_size - размер входных данных в элементах
    __shared__ float avg[AVG2_THREADS_IN_BLOCK];
    
    unsigned int index = blockIdx.x*blockDim.x + threadIdx.x;
    // копируем данные в shared и делим их на data_size
    if (index<data_size) {
        avg[threadIdx.x]=float(d_in_data[index])/data_size;
    }
    else {
        avg[threadIdx.x]=0.;
    }
    __syncthreads();

        // считаем сумму в shared через scan. Результат в последней ячейке
    for (unsigned int i=2; i<=blockDim.x; i<<=1) {
        if ((threadIdx.x +1) % i == 0) {
            avg[threadIdx.x]+=avg[threadIdx.x-i/2];
        }
    }
    __syncthreads();

    if (threadIdx.x==blockDim.x-1) { // последняя задача выдаёт результат
        atomicAdd(d_out_value,float(avg[threadIdx.x]));
    }
}


// функция обёртка над kernel для вычисления среднего арифметического

bool CUDA_Average2(WORD* d_in_data, UINT data_size, float& result) {
        
    // определяем кол-во блоков.
    UINT blocksNum=data_size/AVG2_THREADS_IN_BLOCK;
    if (data_size%AVG2_THREADS_IN_BLOCK!=0) blocksNum++;
    if (blocksNum==0) blocksNum++;

    dim3 threads(AVG2_THREADS_IN_BLOCK);
    dim3 blocks(blocksNum);

    // выделяем память под массив промежуточных результатов
    float* d_out_value=0; float h_out_value=0;
    try {
        if (cudaMalloc(&d_out_value, sizeof(float)) !=  cudaSuccess) throw (1);
        if (cudaMemcpy(d_out_value, &h_out_value, sizeof(float), cudaMemcpyHostToDevice) !=  cudaSuccess) throw (1);
    
        average_kernel_WORDS2<<<blocks,threads,sizeof(float)*AVG2_THREADS_IN_BLOCK>>>(d_in_data,data_size,d_out_value);
        if (cudaMemcpy(&h_out_value, d_out_value, sizeof(float), cudaMemcpyDeviceToHost) !=  cudaSuccess) throw (1);
    }
    catch (...) {
        if (d_out_value!=0) cudaFree(d_out_value);
        return false;
    }

    result=h_out_value;
    cudaFree(d_out_value);

    return true;
}

Bitman

Дата 10.7.2015, 23:14 (ссылка)

(нет голосов)

Загрузка ...

Новичок

Профиль
Группа: Участник
Сообщений: 6
Регистрация: 7.6.2005

Репутация: нет
Всего: нет

Продолжаю копать. Где-то всё равно есть гонки в алгоритме :(
Вот результаты тестовых прогонов версии Release.
Первое число считается serial алгоритмом, сначала идёт суммирование массива 1M в double, а в конце делим на 1M.
Второе число тоже serial, только тут каждое делится на 1M, а потом всё это суммируется.
Третье число - считается на CUDA аналогично 2-му варианту. Явно где-то гонки, но где?!

Цитата

>CudaMedianaTest1.exe
True avg: 1462.14 Time: 0.001248 msec.
Next avg: 1462.14 Time: 0.001216 msec.
GPU Avg: 1234.64 Time: 4.16483 msec.

>CudaMedianaTest1.exe
True avg: 1462.14 Time: 0.001216 msec.
Next avg: 1462.14 Time: 0.001248 msec.
GPU Avg: 1241.22 Time: 4.31347 msec.

>CudaMedianaTest1.exe
True avg: 1462.14 Time: 0.001216 msec.
Next avg: 1462.14 Time: 0.001216 msec.
GPU Avg: 1237.92 Time: 4.03677 msec.

>CudaMedianaTest1.exe
True avg: 1462.14 Time: 0.001216 msec.
Next avg: 1462.14 Time: 0.001248 msec.
GPU Avg: 1235.58 Time: 4.08221 msec.

>CudaMedianaTest1.exe
True avg: 1462.14 Time: 0.001216 msec.
Next avg: 1462.14 Time: 0.001216 msec.
GPU Avg: 1234.84 Time: 4.07472 msec.

Bitman

Дата 12.7.2015, 10:32 (ссылка)

(нет голосов)

Загрузка ...

Новичок

Профиль
Группа: Участник
Сообщений: 6
Регистрация: 7.6.2005

Репутация: нет
Всего: нет

Код


    // считаем сумму в shared через scan. Результат в последней ячейке
    for (unsigned int i=2; i<=blockDim.x; i<<=1) {
        if ((threadIdx.x +1) % i == 0) {
            avg[threadIdx.x]+=avg[threadIdx.x-i/2];
        }
        __syncthreads();
    }

Ларчик просто открывался, нужен был барьер внутри цикла.

0 Пользователей читают эту тему (0 Гостей и 0 Скрытых Пользователей)
0 Пользователей:
« Предыдущая тема \| C/C++: Общие вопросы \| Следующая тема »