CUDA, копирование в разделяемую память резко увеличивает количество регистров, используемых - PullRequest
Новая
       31

CUDA, копирование в разделяемую память резко увеличивает количество регистров, используемых

0 голосов
/ 11 июля 2019

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

Я работаю с графическим процессором Nvidia Tesla K40c. Я использую Pycuda для решения системы PDE. Итак, моя цель - выполнить локальные вычисления для каждого потока, а затем записать их в массив общей памяти. Я довольно плохо знаком с вычислениями на GPU, но то, что я знаю об этой проблеме, написано ниже. Эта проблема связана с закомментированной строкой кода, приведенной ниже. Я знаю, что разделяемая память идеально подходит для межпотоковой связи в блоке, и моя разделяемая память работает правильно, пока я не попытаюсь записать в нее локальную переменную, которая, как я предполагаю, хранится в регистрах. Я предполагаю это, потому что я прочитал, что массивы меньше определенного размера, 16 операций с плавающей запятой, если я правильно помню, МОГУТ храниться в регистрах. Мои имеют размер 4. В любом случае, это цель избежать хранения в глобальном масштабе. 

__device__
void step(float *shared_state, int idx)
{
  float dfdxy[NVC]={0};
  get_dfdx(dfdxy,shared_state,idx);
  get_dfdy(dfdxy,shared_state,idx);
  __syncthreads();
  //shared_state[idx+0*SGIDS] += dfdxy[0];
}

Вот след. Как я уже говорил, я знаком с ошибкой. 

Traceback (most recent call last):
  File "./src/sweep/sweep.py", line 325, in <module>
    sweep(arr0,targs,order,block_size,euler.step,0)
  File "./src/sweep/sweep.py", line 109, in sweep
    gpu_speed(arr, source_mod, cpu_fcn, block_size,ops,num_tries=20)
  File "./src/sweep/sweep.py", line 175, in gpu_speed
    gpu_fcn(arr_gpu,grid=grid_size, block=block_size,shared=shared_size)
  File "/home/walkanth/.conda/envs/pysweep/lib/python3.6/site-packages/pycuda/driver.py", line 402, in function_call
    func._launch_kernel(grid, block, arg_buf, shared, None)
pycuda._driver.LaunchError: cuLaunchKernel failed: too many resources requested for launch

Проблема заключается именно в этом, когда я запускаю код с прокомментированной строкой. Это говорит о том, что я использую 32 регистра. Это хорошо, все работает, потому что я ниже 63.

Однако, когда я раскомментирую строку, количество используемых регистров возрастает до 70, и я подозреваю, что именно поэтому запуск ядра завершается неудачей.

Итак, пара вопросов.

Во-первых, кто-нибудь может объяснить, почему это происходит? Я искал некоторое время и потерпел неудачу.

Во-вторых, если нет способа обойти это. Кто-нибудь знает несколько советов, чтобы уменьшить использование моего регистра помимо уменьшения block_size? Я видел, как некоторые старые темы на nvidia dev говорили об этом, но они, похоже, устарели.

Edit:

Так что, благодаря Майклу в этом посте, я узнал, что у меня на самом деле у GPU 255 регистров на поток. Таким образом, регистры не проблема. Однако из-за этого я не уверен, откуда возникла проблема.

Я также подумал, что было бы полезно добавить, что я не использую какие-либо конкретные параметры компилятора. Я однажды попробовал -ptxas, но он не сильно изменился.

Я не хочу уменьшать размер блока, потому что количество вычислений, которые я могу выполнить до получения внешней информации, зависит от минимального размера блока (x или y). Чем больше размер блока, тем больше расчетов возможно.

Edit: Итак, насколько я понимаю, я все еще превышаю общее количество регистров на SM, что является причиной проблемы. Мне нужно уменьшить количество используемых регистров или размер блока.

1 Ответ

0 голосов
/ 11 июля 2019

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

Однако, согласно https://developer.nvidia.com/cuda-gpus,, K40c имеет вычислительную способность 3.5, ив соответствии с https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capabilities, устройства, которые имеют вычислительные возможности 3.5, могут иметь до 255 регистров на поток, а не 63. Таким образом, если вы все еще используете только 70 регистров на поток, то это, вероятно, не проблема.Это подтверждается, если вы больше не получаете ошибку, уменьшая размер блока;уменьшение размера блока уменьшает количество потоков в блоке, но не должно изменять количество регистров, используемых в каждом потоке, поэтому это не должно решить вашу проблему, если у вас фактически кончились регистры на поток.

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

...