Ограничения на размер grid в Nvidia CUDA при двухмерной сетке?

Question

[al_indigo]

Всем привет. Возможно, мой вопрос покажется многим на первый взгляд очевидным, но я бы всё-таки призвал вас не считать его идиотским до тех пор, пока не прочтёте до конца.


Итак, в чём же суть вопроса. Как известно из документации CUDA, размер сетки запускаемого ядра имеет ограничения, которые зависят от конкретного устройства. В большинстве современных видеокарт ограничение равно 65535x65535x1. На моих видеокартах g210m и 8800gt это именно так, я проверял. Но в этом месте я встретил довольно странную вещь — в моей программе по какой-то неизвестной мне причине невозможно запустить ядро, которое бы имело размерность (по нитям) больше 5808x5808 (это число может быть меньше в зависимости от размера блока, я написал строгий максимум) или же больше 264х264 (если измерять в блоках) — и последнее число неизменно. Как только количество запускаемых блоков переваливает за 265х265, ядро запускается, отрабатывает, но в качестве результата всегда выдаёт ноль.

Дебаггер от Nvidia Nsight молчит, никаких ошибок не выбрасывается, профайлер выдаёт результаты работы, в которых ядро запускается. Ограничение всплывает на всех видеокартах, на которых я запускал программу — в сумме на 8 разных моделях (8400m g, 8800gt, 9600gso, 8500gt, 9600gt, ION, g210m, gf9300)


Так вот всё это наводит меня на мысль, что есть ограничения не только на размерность сетки, но и на суммарное количество нитей в сетке (ведь на количество нитей в блоке ограничение есть — почему бы и тут ему не быть). Только вот ни официальная документация, ни учебник Борескова/Харлмова, ни best practices guide ничего не этот счёт не говорят — просто говорят, что есть ограничения, уже озвученные в самом начале вопроса.


Поскольку копаюсь я с этим примерно по часа два в день на протяжении уже недели, и никакого прогресса нет, я прошу помощи — куда копать? Любые замечания приветствуются, если нужно сделать какие-то уточнения — скажите

Comments

[liq]

Что значит «в качестве результата всегда выдаёт ноль»?

[al_indigo]

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

Answer 1

[liq]

Только что проверил. Мне не удалось повторить вашу проблему.
У меня GTX470.
Итак. Написал ядро:

__global__ void testKernel( int* g_odata) 
{
  if(threadIdx.x==0)
  {
    g_odata[2*(blockIdx.y*gridDim.x+blockIdx.x)] = blockIdx.y;
    g_odata[2*(blockIdx.y*gridDim.x+blockIdx.x)+1] = blockIdx.x;
  }
}

Запустил его на 8192х8192 блоков и 1024 треда(в ваших видяхах максимум 512 тредов в блоке, на ферми 1024):

    dim3  grid( 8192, 8192, 1);
    dim3  threads( 1024, 1, 1);
    testKernel<<< grid, threads, 0 >>>(  d_odata);

Естественно выделил память и т.п.
И получил последним элементом массива: 8191x8191.
На больших числах не тестировал, потому что память кончается :( Надо уже какую-то логику реализовывать.

А вообще не понятно откуда у вас эти не круглые значения 265, 264?

Comments

[al_indigo]

Я в ступоре, на самом деле. Похоже, я где-то всё-таки ошибся, но как-то очень уж странно, что я никак это отловить не могу и что программа почти везде работает правильно
Спасибо, что проверили!

[al_indigo]

А Вы не могли бы попробовать запустить вот нечто такое:

__global__ void testKernel( int* g_odata)
{
int indexX = blockIdx.x*blockDim.x + threadIdx.x;
int indexY = blockIdx.y*blockDim.y + threadIdx.y;
if(indexX==2097088 && indexY==2097088)
{
g_odata[0] = indexX;
g_odata[1] = indexY;
}
__suncthreads();
}

dim3 grid( 65534, 65534, 1);
dim3 threads( 32, 32, 1);
testKernel<<< grid, threads, 0 >>>( d_odata);

[liq]

Попробовал. Конкретно ваш код вылетает по timeout. Подробнее про таймауты тут: forums.nvidia.com/lofiversion/index.php?t106635.html.

Немного модифицировав код получил такой результат:

Checking 1x1
Processing time: 55.926998 (ms)
Last 31x31
—
Checking 2x2
Processing time: 0.098000 (ms)
Last 63x63
—
…
— Checking 256x256
Processing time: 3.470000 (ms)
Last 8191x8191
— …
— Checking 8192x8192
Processing time: 3465.157959 (ms)
Last 262143x262143
— Checking 16384x16384
Processing time: 13827.656250 (ms)
Last 524287x524287
— Checking 32768x32768
template.cu(98): cudaSafeCall() Runtime API error: the launch timed out and was terminated.

[al_indigo]

Спасибо, то, что нужно.
Значит, ошибка где-то у меня, к моему сожалению