Глюк CUDA?

Алексей Вахтин

unread,

Sep 5, 2013, 6:53:33 PM9/5/13

to cudac...@googlegroups.com

Столкнулся с такой проблемой, в результате которой обнаружил следующее недоразумение. Если это не глюк, то объясните, почему оно так работает. Если написать процедуру, где N=5000, M=1024, fromM=1:

__global__ void cuftFurieNew(float* in_data, float* out_data, int N, int M, int fromM = 1)

{

cuComplex s;

float t;

for (int I = threadIdx.x + blockIdx.x * blockDim.x + fromM; I <= M; I += blockDim.x * gridDim.x)

{

s.x = 0;

s.y = 0;

for (int J = 0; J < N; J++)

{

t = (6.0 * (J - N / 2));

// s.x += in_data[J] * cosf(t);

// s.y += in_data[J] * sinf(t);

}

out_data[I - fromM] = I; //s.x * s.x + s.y * s.y;

}

То в out_data попадут числа от 1 до 1024. Если написать

__global__ void cuftFurieNew(float* in_data, float* out_data, int N, int M, int fromM = 1)

{

cuComplex s;

float t;

for (int I = threadIdx.x + blockIdx.x * blockDim.x + fromM; I <= M; I += blockDim.x * gridDim.x)

{

s.x = 0;

s.y = 0;

for (int J = 0; J < N; J++)

{

t = (6.0 * (J - N / 2));

// s.x += in_data[J] * cosf(t);

// s.y += in_data[J] * sinf(t);

}

out_data[I - fromM] = t / 1024; //s.x * s.x + s.y * s.y;

}

все выведется корректно, но если написать:

__global__ void cuftFurieNew(float* in_data, float* out_data, int N, int M, int fromM = 1)

{

cuComplex s;

float t;

for (int I = threadIdx.x + blockIdx.x * blockDim.x + fromM; I <= M; I += blockDim.x * gridDim.x)

{

s.x = 0;

s.y = 0;

for (int J = 0; J < N; J++)

{

t = (6.0 * (J - N / 2));

// s.x += in_data[J] * cosf(t);

// s.y += in_data[J] * sinf(t);

}

out_data[I - fromM] = t / I; //s.x * s.x + s.y * s.y;

}

то в out_data будет мусор, по видимому процедура вываливается с ошибкой. Почему?

Alexander Guzhva

unread,

Sep 5, 2013, 6:55:37 PM9/5/13

to cudac...@googlegroups.com

Например, может произойти деление на 0, если I = 0

5 сентября 2013 г., 18:53 пользователь Алексей Вахтин <alva...@gmail.com> написал:

--
Вы получили это сообщение, поскольку подписаны на группу CUDA.CS.MSU.SU.

Чтобы отказаться от подписки на эту группу и перестать получать из нее сообщения, отправьте электронное письмо на адрес cudacsmsusu...@googlegroups.com.
Чтобы добавлять сообщения в эту группу, отправьте письмо по адресу cudac...@googlegroups.com.
Перейдите в группу по ссылке http://groups.google.com/group/cudacsmsusu.
Настройки подписки и доставки писем: https://groups.google.com/groups/opt_out.

Dmitry N. Mikushin

unread,

Sep 5, 2013, 7:02:02 PM9/5/13

to cudac...@googlegroups.com

Ну если t - это float, то деление на ноль должно дать не ошибку, а Inf. Объяснения лучше искать следующими способами:

0) Всегда проверять код возврата CUDA-процедур и ядра

1) Использовать cuda-memcheck

2) Использовать отладчик
3) Прикладывать к письму полный, но максимально упрощённый пример, демонстрирующий проблему, чтобы мы могли проверить его не только глазами, но и в работе.

6 сентября 2013 г., 0:55 пользователь Alexander Guzhva <alexand...@gmail.com> написал:

Alexander Guzhva

unread,

Sep 5, 2013, 7:04:55 PM9/5/13

to cudac...@googlegroups.com

Исхожу из того, что автор все это проверил. Могу предположить только больной случай, что CUDA компилятор сильно умный, увидел, что значения, присваемые t - только целые, и делит t / I как int на int.

5 сентября 2013 г., 19:02 пользователь Dmitry N. Mikushin <maem...@gmail.com> написал:

Dmitry N. Mikushin

unread,

Sep 5, 2013, 7:32:30 PM9/5/13

to cudac...@googlegroups.com

Можно представить себе сочетание условий:

1) Если пользователь не проверяет код возврата и диапазоны массивов, то ядро может быть некорректным и при этом получать [проверяемую] часть правильного результата

2) Порядок обработки блоков неопределён, но известно что он (а) примерно неизменен при повторном запуске одного и того же кода, однако (б) может как угодно меняться при изменении кода. За счёт этого, в случае GPU (или динамической многопоточности на CPU) упрощение программы может давать потенциально более различный control flow, чем в случае однопоточного кода, где - действительно - только фактор компилятора.

- Д.

6 сентября 2013 г., 1:04 пользователь Alexander Guzhva <alexand...@gmail.com> написал:

Алексей Вахтин

unread,

Sep 6, 2013, 3:08:01 AM9/6/13

to CUDA.CS.MSU.SU

то в out_data будет мусор, по видимому процедура вываливается с ошибкой. Почему?

Я стал думать почему в этом случае не работает, хотя, в других случаях работало и обнаружил следующее. Для того, чтобы использовать atomicAdd мне пришлось изменить параметры компилятора. А именно, в свойствах проекта в разделе CUDA C/C++ -> Device -> Code Generation задал compute_20,sm_20 - первое попавшееся решение на ошибку atomic is undefined. Когда я вернул прежние значения - compute_10,sm_10 все стало работать. Еще я выяснил, что atomicAdd работает и с compute_11,sm_11 и у меня все работает без проблем.

Как я понимаю, compute и sm задают архитектуру CUDA. Вот только не понятно, это особенности архитектуры 2.0 или видеокарта не поддерживает эту архитектуру? Видеокарта MSI GeForce GTX 550 Ti, операционная система - Windows 8 x64.

Alexey V. Medvedev

unread,

Sep 6, 2013, 5:36:15 AM9/6/13

to cudac...@googlegroups.com

Алексей,

compute_20,sm_20 -- несомненно поддерживаются этой видеокартой, оставьте их и не меняйте обратно! Всё должно работать, надо искать ошибку в коде.

Как уже говорили, надо проследить, чтобы везде проверялись коды ошибок, после каждого без исключения вызова API и после каждого вызова ядра (через cudaDeviceSynchronize() например), а также собрать код с отладочными ключами (-O0 -G), запустить под nsight обязательно с включённым memory checker. Без этого дальнейший анализ не имеет смысла.

--

Regards,

Alexey

2013/9/6 Алексей Вахтин <alva...@gmail.com>

то в out_data будет мусор, по видимому процедура вываливается с ошибкой. Почему?

Я стал думать почему в этом случае не работает, хотя, в других случаях работало и обнаружил следующее. Для того, чтобы использовать atomicAdd мне пришлось изменить параметры компилятора. А именно, в свойствах проекта в разделе CUDA C/C++ -> Device -> Code Generation задал compute_20,sm_20 - первое попавшееся решение на ошибку atomic is undefined. Когда я вернул прежние значения - compute_10,sm_10 все стало работать. Еще я выяснил, что atomicAdd работает и с compute_11,sm_11 и у меня все работает без проблем.

Как я понимаю, compute и sm задают архитектуру CUDA. Вот только не понятно, это особенности архитектуры 2.0 или видеокарта не поддерживает эту архитектуру? Видеокарта MSI GeForce GTX 550 Ti, операционная система - Windows 8 x64.

--

Алексей Вахтин

unread,

Sep 8, 2013, 4:28:27 PM9/8/13

to CUDA.CS.MSU.SU

Спасибо откликнувшимся. Пытался проанализировать обнаруженную мной проблему. Ошибок не обнаружил - всегда возвращается cudaSuccess. Как работать с NSight не знаю (дайте, пожалуйста, ссылку на мануал как запустить отладку обычной программы). Вот упрощенный код программы, выбросил по возможности все лишнее:

#include <cufft.h>

#include <stdio.h>

#include "cuda.h"

#include "cuda_runtime.h"

#include "device_launch_parameters.h"

#pragma once //чтобы скрыть ошибки во время редактирования

#ifdef __INTELLISENSE__

void __syncthreads();

void atomicAdd(int*, int);

#endif

__global__ void cuftFurieNew(float* in_data, float* out_data, int N, int M, int fromM = 1)

{

cuComplex s;

float t = 0;

for (int I = threadIdx.x + blockIdx.x * blockDim.x + fromM; I <= M; I += blockDim.x * gridDim.x)

{

s.x = 0;

s.y = 0;

for (int J = 0; J < N; J++)

{

t = (6.0 * (J - N / 2)) / I;

s.x = in_data[J] * cos(t);

s.y = in_data[J] * sin(t);

}

out_data[I - fromM] = 400;//s.x * s.x + s.y * s.y;

}

extern "C" __declspec(dllexport) void cuftColorQualifierExec(float* data, float *spm, int N, int M, int fromM)

{

float* in_data_dev;//начальные данные на устройстве

float *furie_dev;//выходные данные на устройстве - сам вейвлет, результат обратного ПФ

//инициализация работы с кудой

cudaDeviceProp prop;//свойства устройства

int N_Dev;//номер устройства

memset(&prop, 0, sizeof(cudaDeviceProp));//обнуляем

prop.major = 2;//задается версия куды

prop.minor = 0;

prop.maxThreadsPerBlock = M - fromM;//задаем желаемое количество нитей в блоке

cudaChooseDevice(&N_Dev, &prop);//получаем номер наиболее подходящее устройство

cudaSetDevice(N_Dev);//задаем устройство

cudaGetDeviceProperties(&prop, N_Dev);//получаем свойства устройства

//вычисляем количестов нитей и блоков в соответствии с объемом данных и максимальными возможностями устройстве*/

(************************ если здесь поменять, то будет работать ****************************)

int N_thread = min(prop.maxThreadsPerBlock, prop.maxThreadsDim[0]);

int N_block = 1;//prop.maxGridSize[0];

(******************************************************************************************************)

int *Count_dev;//число максимумов

//выделяем память для начальных данных

cudaError_t err = cudaMalloc((void**)&in_data_dev, sizeof(float) * N);

if (err != cudaSuccess)