Функция CUDA __umul24, полезна или нет? - PullRequest
Новая
       28

Функция CUDA __umul24, полезна или нет?

4 голосов
/ 05 апреля 2011

Стоит ли заменять все умножения функцией __umul24 в ядре CUDA?Я читаю разные и противоположные мнения, и я до сих пор не могу сделать отметку, чтобы понять это

Ответы [ 2 ]

3 голосов
/ 11 декабря 2011

Просто хотел присоединиться к нему с несколько иным мнением, чем Эшвин / ФабрициоМ ...

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

Скорее всего, многие пользователи, которые будут использовать CUDA на устаревших компьютерах, получат преимущества от использования соответствующих функций уровня вычислений. И это не так сложно, как это делает Ashwin / fabrizioM.

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

//For prior to Fermi use umul, for Fermi on, use
//native mult.
__device__ inline void MultiplyFermi(unsigned int a, unsigned int b)
{ a*b; }

__device__ inline void MultiplyAddFermi(unsigned int a, unsigned int b,
                                        unsigned int c)
{ a*b+c; }

__device__ inline void MultiplyOld(unsigned int a, unsigned int b)
{ __umul24(a,b); }

__device__ inline void MultiplyAddOld(unsigned int a, unsigned int b,
                                      unsigned int c)
{ __umul24(a,b)+c; }

//Maximum Occupancy =
//16384
void GetComputeCharacteristics(ComputeCapabilityLimits_t MyCapability)
{
    cudaDeviceProp DeviceProperties;
    cudaGetDeviceProperties(&DeviceProperties, 0 );
    MyCapability.ComputeCapability =
    double(DeviceProperties.major)+ double(DeviceProperties.minor)*0.1;
}

Теперь здесь есть обратная сторона. Что это?

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

Стоит ли это того?

Хорошо, подумайте, это тривиальная работа по копированию и вставке, и вы набираете эффективность, да, по моему мнению. В конце концов, CUDA не является самой легкой концептуальной формой программирования (как и любое параллельное программирование). Если производительность НЕ критична, спросите себя: почему вы используете CUDA?

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

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

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

3 голосов
/ 05 апреля 2011

Только в устройствах с архитектурой до Fermi, то есть с возможностями cuda до 2.0, где целочисленная арифметическая единица составляет 24 бита.

На устройстве Cuda с возможностями> = 2.0 архитектура 32-битная, _umul24 будет медленнее, чем быстрее. Причина в том, что он должен эмулировать 24-битную операцию с 32-битной архитектурой.

Вопрос теперь: стоит ли усилий для увеличения скорости? Вероятно, нет.

...