Глюк CUDA?

186 views
Skip to first unread message

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

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)
fprintf(stderr, "ERROR \"%s:%d\n", cudaGetErrorString(err), __FILE__, __LINE__);
//копируем данные в память
err = cudaMemcpy(in_data_dev, data, sizeof(float) * N, cudaMemcpyHostToDevice);
if (err != cudaSuccess)
fprintf(stderr, "ERROR \"%s:%d\n", cudaGetErrorString(err), __FILE__, __LINE__);
//выделяем память для выходных данных
err = cudaMalloc((void**)&furie_dev, sizeof(float) * (M - fromM + 1));
if (err != cudaSuccess)
fprintf(stderr, "ERROR \"%s:%d\n", cudaGetErrorString(err), __FILE__, __LINE__);
//N_thread = 1; N_block = 5651;
//запускаем процесс
cuftFurieNew<<<N_block, N_thread>>>(in_data_dev, furie_dev, N, M, fromM); 

err = cudaDeviceSynchronize();
if (err != cudaSuccess)
fprintf(stderr, "ERROR \"%s:%d\n", cudaGetErrorString(err), __FILE__, __LINE__);

//копируем spm
err = cudaMemcpy(spm, furie_dev, sizeof(float) * (M - fromM + 1), cudaMemcpyDeviceToHost);
if (err != cudaSuccess)
fprintf(stderr, "ERROR \"%s:%d\n", cudaGetErrorString(err), __FILE__, __LINE__);

//Удаляем spm с устройства
err = cudaFree(furie_dev);
if (err != cudaSuccess)
fprintf(stderr, "ERROR \"%s:%d\n", cudaGetErrorString(err), __FILE__, __LINE__);
}

int main()
{
int M = 1024, fromM = 1, N = 5000;
float* data = new float[5000];
float* spm = new float[M - fromM + 1];

for (int I = 0; I < N; I++)
data[I] = cos(6.0 * I / 2) + cos(6.0 * I / 3) + cos(6.0 * I / 4) + cos(6.0 * I / 5) + cos(6.0 * I / 6) + cos(6.0 * I / 7);

for (int I = 0; I < M - fromM + 1; I++)
spm[I] = 0;

cuftColorQualifierExec(data, spm, N, M, fromM);

for (int I = 0; I < M - fromM + 1; I++)
fprintf(stdout, "%d: %f\n", I, spm[I]);

    return 0;
}

Жирным я выделил то, что если заблокировать, то в элемент массива попадет число 400. Иначе - нет. Можно поменять число, чтобы убедиться, работает или нет, так как иногда после успешной работы остается в памяти старое значение. Жирным курсивом выделил то, что если поменять, то все будет работать. А именно - число нитей и блоков. Если оно задано:

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

То не работает. А если задать:

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

То будет работать. Не могу понять, почему? Память выделяется правильно, деление на 0 исключено, выход за пределы массива - тоже. Достаточно посмотреть на код программы, очевидно. Длина созданного массива, куда копируется не превышает M-fromM+1, а I проходит от fromM до M.

И еще, если не трудно, дайте ссылку, где можно прочитать про нити и блоки. Еще одну странную вещь обнаружил. Если скомпилировать с параметрами compute_11,sm_11 и задать:

int N_thread = 1;
int N_block = 5651;

То после работы ядра 

        cuftFurieNew<<<N_block, N_thread>>>(in_data_dev, furie_dev, N, M, fromM); 

возвращается ошибка unknownerror, при этом windows сообщает, что произошел сбой драйвера. При этом в этом случае программа не выдает ошибок если заблокировать sin и cos. При других значениях параметров N_thread и N_block ошибок не возникает. 



С уважением, Вахтин А. А.

Dmitry N. Mikushin

unread,
Sep 9, 2013, 8:02:52 AM9/9/13
to cudac...@googlegroups.com
Алексей, у меня Ваш код работает (GTX680M, Linux) и с N_block = 1, и с N_block = 2. Пара исправлений:

1) В вызове fprintf(stderr, "ERROR \"%s:%d\n", cudaGetErrorString(err), __FILE__, __LINE__) не хватает аргумента. Заменил на макрос
2) Размер грида на sm_30 и выше больше, чем диапазон int. Заменил на uint64_t, и, соответственно, тип индекса в ядре. Иначе int переполнится, станет отрицательным, и всё пропало.

Но это скорее косметика. А вот принципиальная проблема состоит в том, что автор кода не до конца разобрался в том, как выбирать конфигурацию грида. В частности:

1) Возможность или невозможность запуска конфигурации ограничена не только размерностями, но и, например, числом регистров. Почитайте об этом в книге или в доках. Скорее всего работающий/не работающий случаи объясняются именно этим.
2) Непонятно, какой смысл запускать ядро с гридом, не зависящим от размера самой задачи. Скажем вот с uint64_t N_block = prop.maxGridSize[0] будет честно запущен миллиард блоков. На свободном GPU это будет считаться очень долго, а на GPU с таймаутом драйвер остановит ядро через небольшое время и выдаст ошибку (как раз упоминаемый Вами "сбой драйвера").

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

- Д.



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

--
vakhtin.cu
Reply all
Reply to author
Forward
0 new messages