CUDA на примерах

Подозреваю, что ты гвоздями вбил ограничения: на количество потоков, на размер памяти и т.п.

А еще есть такое узкое место, как копирование между оперативкой и картой. Возможно, что у тебя затык из-за этого.

Попробуй что-нибудь тяжелое: например, Фурье-фильтрацию изображения. Если все равно время выполнения будет одинаковым, что-то тут нечисто.

Может это быть из-за драйвера ASUS? Может установить драйвер с сайта nvidia?

Это уже welcome to winFAQ

После меняю карту на новую: NVidia GeForce GT 630. Все - то же самое, но время выполнения в данных вариантах - практически одно и тоже, вне зависимости от памяти, при чем не самое маленькое.

Хоть бы сказал, что за алгоритм. А так, не исключено кеширование глобальной памяти, запиленное в Fermi: оно может несколько нивелировать латентность глобальной памяти.

Нопочему на старой заведомо худшей карте все работало как надо?..Вот в чем вопрос...

Где код?

Код из книги: программа трасссировщика лучей.Привожу вариант с текстурной памятью. void anim_gpu( DataBlock *d, int ticks ) { cudaEventRecord( d->start, 0 ); dim3 blocks(DIM/16,DIM/16); dim3 threads(16,16); CPUAnimBitmap *bitmap = d->bitmap; // since tex is global and bound, we have to use a flag to // select which is in/out per iteration volatile bool dstOut = true; for (int i=0; i<90; i++) { float *in, *out; if (dstOut) { in = d->dev_inSrc; out = d->dev_outSrc; } else { out = d->dev_inSrc; in = d->dev_outSrc; } copy_const_kernel<<<blocks,threads>>>( in ); blend_kernel<<<blocks,threads>>>( out, dstOut ); dstOut = !dstOut; } float_to_color<<<blocks,threads>>>( d->output_bitmap, d->dev_inSrc ); cudaMemcpy( bitmap->get_ptr(), d->output_bitmap, bitmap->image_size(), cudaMemcpyDeviceToHost );

cudaEventRecord( d->stop, 0 ); cudaEventSynchronize( d->stop ); float elapsedTime; cudaEventElapsedTime( &elapsedTime, d->start, d->stop );

d->totalTime += elapsedTime; ++d->frames; printf( «Average Time per frame: %3.1f ms\n», d->totalTime/d->frames ); } Думаю,понятно...

Почитай про LORcode

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

LORcode - что это?

www.linux.org.ru/wiki/en/Lorcode

Только текстурная и константная памяти могут иметь кэш,насколько я понимаю...

dim3 blocks(DIM/16,DIM/16); 
dim3 threads(16,16);

Неужели это самые оптимальные размеры?

CUDA только изучаю, потому примеры - из книги.

Говорят же, в Fermi сделали L1 кэш для глобальной памяти.

Вопрос - почему это «гвоздями прибитое количество потоков, и активное копирование данных» - работало на старой карте как надо (со временем)?..

И да, советую почитать

• David B.Kirk, Wen-mei W.Hwu — Programming Massively Parallel Processors A Hands-on Approach
• Боресков А.В., Харламов А.А. - основы работы с технологией CUDA

В этих книгах больше внимания уделено вопросам оптимизации.

Пересобери с –Xptxas –dlcm=cg. Если версия без текстурной и shared памяти начнёт работать медленнее, значит дело в кэше.

То есть книга «CUDA на примерах» уже устарела...

А потому что на старой карте было мало ядер. И мало памяти. И все было ОК. А на новой ядер много, но куча простаивает из-за неправильного распределения ресурсов.

кстати, сейчас идёт курс этого wen-mei W Hwu - https://www.coursera.org/course/hetero

ещё не поздно подключиться. думаю.