Как сделать модуль комплексного числа более эффективным в CUDA?

Question

Я пытался выполнить быстрое преобразование Фурье для собранных данных. После операции БПФ я хотел вычислить модуль данных типа cufftComplex. Поэтому я суммировал квадрат действительной части и квадрат мнимой части, а затем взял квадрат root суммирования. Ниже приведен код, а также назначение сеток и блоков:

dim3 dimBlock(256);
dim3 dimGrid(FFTlength / 256 * lines);

__global__ void modulus_kernel(int length, int lines, cufftComplex *PostFFTData, float* z) 
{
    unsigned int x = blockIdx.x * blockDim.x + threadIdx.x;
    if(x<length*lines)
        z1[x] = sqrt(PostFFTData[x].x *PostFFTData[x].x + PostFFTData[x].y *PostFFTData[x].y);
    __syncthreads();
}

Длина массива указателей PostFFData составляет 1024000, а длина и строки - 2048 и 500 соответственно. После того, как я выполнил код, я проанализировал временную шкалу программы с помощью Nvidia Visual Profiler. Это показывает, что для завершения модуля ядра потребовалось 0,367 мс. Кроме того, я использовал видеокарту GTX1080, а CPU - i7-7700U. Если я хочу сократить время выполнения, как мне это сделать?

Answer 1

Если я хочу сократить время выполнения, как мне это сделать?

Я могу придумать как минимум пять вещей (в произвольном порядке)

1. Избавьтесь от звонка __syncthreads(). В этом нет необходимости и он будет активно замедлять ваш код.
2. Передайте ядро ​​length*lines как единственный аргумент для ядра. Почему каждый поток выполняет целочисленное умножение на постоянное значение?
3. Используйте шаг сетки l oop и запускайте столько потоков, сколько может находиться на устройстве. Используйте API-интерфейсы занятости , чтобы среда выполнения тщательно обдумывала параметры запуска за вас.
4. Если размер проблемы позволяет использовать #pragma unroll с предлагаемой длиной развертывания, чтобы указать компилятору что размер решетки l oop можно частично раскатывать. Если это не позволяет компилятору генерировать поток операций с плавающей запятой, тогда частично разверните сетку размером l oop самостоятельно.
5. Поскольку вы передаете значения с плавающей запятой одинарной точности, используйте sqrtf, не sqrt. Существуют значительные различия в производительности между функциями двойной и одинарной точности. Если ваше приложение позволяет это, рассмотрите возможность использования менее точных версий функции sqrt (prec-sqrt=false)

Answer 2

__syncthreads();

бесполезно, так как нет разделения между потоками