2015-08-04 6 views
0

Это будет немного напуганный вопрос, который я предполагаю, и если мне нужно уточнить, скажите об этом.Как совместить __restrict__ с массивом, на который указывает указатель __constant__?

Ситуация такова: у меня около 2 гигабайт памяти GPU, содержащих мои случайные числа, и мне нужно использовать их во многих разных функциях. Чтобы предотвратить прохождение вокруг указателей в эту память, от функции устройства до функции устройства (и это много раз), я помещаю указатели в постоянную память gpu, что также сохраняет мне регистры (для меня очень важно). Теперь я знаю, что функции могут ускоряться в некоторых случаях, если они объясняются тем, что фрагменты памяти, на которые указывают его аргументы, не перекрываются, используя ключевое слово __restrict__.

Вопрос: как я могу убедиться, что компилятор знает, что куски памяти в глобальной памяти, на которые указывают указатели в постоянной памяти, не перекрываются (и, возможно, также приятно знать: никогда не изменяются после генерации ядра randoms вызов)?

+2

Как помещать указатели в регистры сохранения постоянной памяти? И вы понимаете, что на всех поддерживаемых в настоящее время архитектурах аргументы ядра передаются в постоянной памяти в любом случае, поэтому использование постоянной памяти для «предотвращения прохода вокруг указателей» не делает ничего, кроме использования обычного аргумента аргумента аргумента, за исключением исключения возможности использования '__restrict__'. – talonmies

+0

Аргументы не только для вызова ядра, но и для многих функций, используемых внутри ядер. Таким образом, от функции устройства до функции устройства мы проходили вокруг этих указателей. Я предположил, что вложение функций уменьшит количество регистров, но я узнал, что, уменьшив количество аргументов, я действительно улучшил число. Вы, однако, ответили на мой вопрос, я думаю: ответ в том, что я не могу, не так ли? Если да, пожалуйста, сделайте это ответом, и я могу согласиться. Я отредактировал мой вопрос, чтобы быть более четким, спасибо за ваш комментарий и потенциальный ответ. – ikku100

+0

Компилятор будет по умолчанию расширять функции '__device__', поэтому, если вы явно не принудительно используете ABI для вызова функций устройства, списки аргументов в функции устройства устраняются компилятором. Это означает, что стандартный аргумент ядра, передаваемый с '__restrict__', будет делать то, что вы хотите, не делая ничего. В противном случае я не знаю, как обеспечить эвристику времени компиляции в других анонимных указателях. – talonmies

ответ

0

Мне неизвестно, как обеспечить компилятор эвристикой в ​​других анонимных указателях.

Если вы можете управлять им, самый простой способ попытаться помочь компилятору выполнить его задачу - передать указатели так, как __restrict__ декорировал аргументы ядра, а затем принудительно включил функции устройства. Это будет обходить ABI и может позволить компилятору использовать известное условие без сглаживания для оптимизации паттернов доступа к памяти. Это также должно помочь с размером регистра данных ваших функций. Я не уверен, что __restrict__ окажет большое влияние на функции __device__ или __constant__, но вы отметили, что компилятор принимает его, поэтому я думаю, что это не помешает, по крайней мере, попробовать.

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