CUDA быстрые математические операции

Question

Это мой код,

    __device__ void calculateDT(float *devD, int *devImg, int cntVoxelLi, int *neighVoxels)
   {
    float minV = devD[cntVoxelLi];
   int cv = devImg[cntVoxelLi];
   float v = 0,cuVal = 0;
   int c1=0,d1=0,r1=0;
   GetInd2Sub(cntVoxelLi, r1,c1,d1);

   for(int ind=0;ind<9;ind++)
   {
    v = pow(float(cv - devImg[neighVoxels[ind]]),2);
    cuVal = devD[neighVoxels[ind]]  + (1-exp(-v/100));
    minV = min(minV, cuVal);
   }
   devD[cntVoxelLi] = minV;
   }

Когда я запускаю всю программу, это занимает около 15 секунд. Но когда я удаляю

     exp(-v/100)

это займет всего 7 секунд. Вроде бы эта операция exp занимает много времени. Я пытаюсь с помощью функции expf. Как я могу улучшить производительность?

Answer 1

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

Что касается оптимизации производительности, очевидным является исключение использования pow для вычисления простого квадрата (компилятор может делать это сам) и приведение в порядок всех выражений с плавающей запятой, чтобы исключить ряд неявных преобразования целых чисел с плавающей точкой (подсказка: 0 - целое число, 0. - двойная точность и 0.f - одинарная точность).

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