Помогите разобраться

Здрасте,
я новичок в этом деле и столкнулся с тем, чего не понимаю:

Есть такой код:

  1. __device__ void initGenerator(unsigned int seed, unsigned int* state_tmp)
  2. {
  3.         state_tmp[0] = seed & 0xffffffffU;
  4.         for(int j=1; j<N; j++)
  5.         {
  6.                 state_tmp[j] = (1812433253U * (state_tmp[j-1] ^ (state_tmp[j-1] >> 30)) + j);
  7.         }
  8. }

N = 32.
Так вот при вызове ЭТОЙ функции видеокарта считает глобальную за 3 000 мс
При вызове такой же функции, но организованной так:

  1. __device__ void initGenerator(unsigned int seed, unsigned int* state_tmp)
  2. {
  3.         state_tmp[0] = seed & 0xffffffffU;
  4.         for(int j=1; j<17; j++)
  5.         {
  6.                 state_tmp[j] = (1812433253U * (state_tmp[j-1] ^ (state_tmp[j-1] >> 30)) + j);
  7.         }
  8.         for(int j=17; j<N; j++)
  9.         {
  10.                 state_tmp[j] = (1812433253U * (state_tmp[j-1] ^ (state_tmp[j-1] >> 30)) + j);
  11.         }
  12. }

глобальная считается за 7 мс !!!!!!!!!!!!! разница в 420 раз! Как ТАК?????

Немного о глобальной:
64 блока, 64 потока в блоке
каждый поток имеет 32 + 3 + 4 около 40 интов. Я так полагаю - пишутся они в регистры.
Разделяемую память не использую, к глобальной вообще не обращаюсь пока не пройдет расчет. Все считается только регистрами (я так надеюсь, контролировать не могу)
Что же влияет на скорость расчета так сильно?
Видяха - GTX 260

зы. результаты для всех потоков одинаковы в обоих случаях... я в замешательстве Т_Т

Forums: 

state_tmp - это ведь у вас в

state_tmp - это ведь у вас в глобальной памяти, правильно? И все потоки ходят в одну и ту же ячейку (state_tmp[j-1]) ?

Вообще, то что 4096 потоков могут делать по ~160 операций (5 в цикле - на глазок) за 3 секунды - это чудо какое-то. 7 миллисекунд - тоже очень много. Вы время как меряете?

ненене state_tmp - это массив

ненене
state_tmp - это массив из 32 интов, такой имеет каждый поток, и я искренне надеюсь, что размещаются они в регистрах (у меня получается 256 регистров на каждый поток в блоке, я расчитывал на выполнение 4 блоков на мультипроцессоре - т.е. 64 регистра на поток).

Приведенная функция в глобальной выполняется 100 000 раз для каждого потока (я это делал для проверки скорости и работоспособности). Вот и не понимаю, как так, почему такая большая разница во времени?

Дальше столкнулся с проблемой по-забавней:
Имеется похожая функция, только там в цикле выполняется поболе действий, так вот ее пришлось разбивать на 4!!!! части (нашел опытным путем)... А самое интересное, если я разбиваю ее руками, как сделал в первом посте (от 0 до 8, от 8 до 16 и тп) - все просто Летает. Если же разделяю программно (цикл в цикле) - картина не меняется, считается очень медленно

Понять бы, чего там не так, если это зависит от видеокарты - труба, мне нужен код для разных моделей

зы. к глобальной памяти я вообще не обращаюсь, только в самом конце, при этом из каждого потока мне нужен только один инт.
ззы. время замерял стандартно

  1. unsigned int timer;
  2. CUT_SAFE_CALL(cutCreateTimer(&timer));
  3. CUT_SAFE_CALL(cutStartTimer(timer));
  4. Kerenl<<< BLOCKS_COUNT, BLOCK_SIZE>>> (trataatatata);
  5. CUT_SAFE_CALL(cutStopTimer(timer));
  6. printf("\nProcessing time GPU: %f (ms) \n", cutGetTimerValue(timer));
  7. CUT_SAFE_CALL(cutDeleteTimer(timer));

Про время - если вы явно не

Про время - если вы явно не зовете синхронизцию после вызова kernel, то время у вас измеряется какое-то. Посмотрите, вот тут вот обсуждали (и не только тут): http://www.gpgpu.ru/node/209

Вообще же, если у вас разбиение цикла на куски с известным числом повторов помогает - единственная мысль, это что компилятор делает loop unroll. Посмотрите код (PTX) в разных случаях. Других идей нет.

я не вписал синхронизацию,

я не вписал синхронизацию, она есть =)

Спасибо за ответы, завтра буду запускать на 280 и 285, проверять работоспособность, может там чего интересного всплывет