Поэлементное векторное умножение с CUDA

Я собрал элементарное ядро ​​в CUDA для поэлементного умножения вектора на вектор двух комплексных векторов. Код ядра вставлен ниже (multiplyElementwise). Это прекрасно работает, но так как я заметил, что другие, казалось бы, простые операции (например, масштабирование вектора) оптимизированы в библиотеках, таких как CUBLAS или CULA, мне было интересно, можно ли заменить мой код вызовом библиотеки? К моему удивлению, ни CUBLAS, ни CULA не имеют этой опции, я пытался подделать ее, сделав один из векторов диагональю диагонального матрично-векторного произведения, но результат был действительно медленным.

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

Итак, мои вопросы:

  1. Есть ли библиотека, которая делает поэлементное вектор-векторное умножение?
  2. Если нет, могу ли я ускорить мой код (multiplyElementwise)?

Любая помощь будет принята с благодарностью!

__global__ void multiplyElementwise(cufftComplex* f0, cufftComplex* f1, int size)
{
    const int i = blockIdx.x*blockDim.x + threadIdx.x;
    if (i < size)
    {
        float a, b, c, d;
        a = f0[i].x; 
        b = f0[i].y;
        c = f1[i].x; 
        d = f1[i].y;
        float k;
        k = a * (c + d);
        d =  d * (a + b);
        c =  c * (b - a);
        f0[i].x = k - d;
        f0[i].y = k + c;
    }
}

__global__ void multiplyElementwiseFast(cufftComplex* f0, cufftComplex* f1, int size)
{
    const int i = blockIdx.x*blockDim.x + threadIdx.x;
    if (i < 4*size)
    {
        const int N = 256;
        const int thId = threadIdx.x / 4;
        const int rem4 = threadIdx.x % 4;
        const int i4 = i / 4;

        __shared__ float a[N];
        __shared__ float b[N];
        __shared__ float c[N];
        __shared__ float d[N];
        __shared__ float Re[N];
        __shared__ float Im[N];

        if (rem4 == 0)
        {
            a[thId] = f0[i4].x;
            Re[thId] = 0.f;
        }
        if (rem4 == 1)
        {
            b[thId] = f0[i4].y;
            Im[thId] = 0.f;
        }
        if (rem4 == 2)
            c[thId] = f1[i4].x;
        if (rem4 == 0)
            d[thId] = f1[i4].y;
        __syncthreads();

        if (rem4 == 0)
            atomicAdd(&(Re[thId]), a[thId]*c[thId]);        
        if (rem4 == 1)
            atomicAdd(&(Re[thId]), -b[thId]*d[thId]);
        if (rem4 == 2)
            atomicAdd(&(Im[thId]), b[thId]*c[thId]);
        if (rem4 == 3)
            atomicAdd(&(Im[thId]), a[thId]*d[thId]);
        __syncthreads();

        if (rem4 == 0)
            f0[i4].x = Re[thId];
        if (rem4 == 1)
            f0[i4].y = Im[thId];
    }
}
Источник

2 ответа

Решение

Если то, что вы пытаетесь достичь, является простым поэлементным продуктом с комплексными числами, вы, похоже, делаете некоторые дополнительные шаги в своем multiplyElementwise ядро, которое увеличивает использование регистра. То, что вы пытаетесь вычислить, это:

f0[i].x = a*c - b*d;
f0[i].y = a*d + b*c;

поскольку (a + ib)*(c + id) = (a*c - b*d) + i(a*d + b*c), Используя улучшенное комплексное умножение, вы торгуете 1 умножением за 3 добавления и несколько дополнительных регистров. Может ли это быть оправдано или нет, может зависеть от используемого вами оборудования. Например, если ваше оборудование поддерживает FMA (Fused Multiply-Add), такая оптимизация может оказаться неэффективной. Вам следует рассмотреть возможность прочтения этого документа: " Точность и производительность: с плавающей точкой и совместимостью IEEE 754 для графических процессоров NVIDIA", в которой также рассматривается проблема точности с плавающей точкой.

Тем не менее, вы должны рассмотреть возможность использования Thrust. Эта библиотека предлагает множество высокоуровневых инструментов для работы с векторами хоста и устройства. Вы можете увидеть длинный список примеров здесь: https://github.com/thrust/thrust/tree/master/examples. Это сделает вашу жизнь намного проще.

ОБНОВЛЕННЫЙ КОД

В вашем случае вы можете использовать этот пример и адаптировать его к чему-то вроде этого:

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <time.h>

struct ElementWiseProductBasic : public thrust::binary_function<float2,float2,float2>
{
    __host__ __device__
    float2 operator()(const float2& v1, const float2& v2) const
    {
        float2 res;
        res.x = v1.x * v2.x - v1.y * v2.y;
        res.y = v1.x * v2.y + v1.y * v2.x;
        return res;
    }
};

/**
 * See: http://www.embedded.com/design/embedded/4007256/Digital-Signal-Processing-Tricks--Fast-multiplication-of-complex-numbers%5D
 */
struct ElementWiseProductModified : public thrust::binary_function<float2,float2,float2>
{
    __host__ __device__
    float2 operator()(const float2& v1, const float2& v2) const
    {
        float2 res;
        float a, b, c, d, k;
        a = v1.x;
        b = v1.y;
        c = v2.x;
        d = v2.y;
        k = a * (c + d);
        d =  d * (a + b);
        c =  c * (b - a);
        res.x = k -d;
        res.y = k + c;
        return res;
    }
};

int get_random_int(int min, int max)
{
    return min + (rand() % (int)(max - min + 1));
}

thrust::host_vector<float2> init_vector(const size_t N)
{
    thrust::host_vector<float2> temp(N);
    for(size_t i = 0; i < N; i++)
    {
        temp[i].x = get_random_int(0, 10);
        temp[i].y = get_random_int(0, 10);
    }
    return temp;
}

int main(void)
{
    const size_t N = 100000;
    const bool compute_basic_product    = true;
    const bool compute_modified_product = true;

    srand(time(NULL));

    thrust::host_vector<float2>   h_A = init_vector(N);
    thrust::host_vector<float2>   h_B = init_vector(N);
    thrust::device_vector<float2> d_A = h_A;
    thrust::device_vector<float2> d_B = h_B;

    thrust::host_vector<float2> h_result(N);
    thrust::host_vector<float2> h_result_modified(N);

    if (compute_basic_product)
    {
        thrust::device_vector<float2> d_result(N);

        thrust::transform(d_A.begin(), d_A.end(),
                          d_B.begin(), d_result.begin(),
                          ElementWiseProductBasic());
        h_result = d_result;
    }

    if (compute_modified_product)
    {
        thrust::device_vector<float2> d_result_modified(N);

        thrust::transform(d_A.begin(), d_A.end(),
                          d_B.begin(), d_result_modified.begin(),
                          ElementWiseProductModified());
        h_result_modified = d_result_modified;
    }

    std::cout << std::fixed;
    for (size_t i = 0; i < 4; i++)
    {
        float2 a = h_A[i];
        float2 b = h_B[i];

        std::cout << "(" << a.x << "," << a.y << ")";
        std::cout << " * ";
        std::cout << "(" << b.x << "," << b.y << ")";

        if (compute_basic_product)
        {
            float2 prod = h_result[i];
            std::cout << " = ";
            std::cout << "(" << prod.x << "," << prod.y << ")";
        }

        if (compute_modified_product)
        {
            float2 prod_modified = h_result_modified[i];
            std::cout << " = ";
            std::cout << "(" << prod_modified.x << "," << prod_modified.y << ")";
        }
        std::cout << std::endl;
    }   

    return 0;
}

Это возвращает:

(6.000000,5.000000)  * (0.000000,1.000000)  = (-5.000000,6.000000)
(3.000000,2.000000)  * (0.000000,4.000000)  = (-8.000000,12.000000)
(2.000000,10.000000) * (10.000000,4.000000) = (-20.000000,108.000000)
(4.000000,8.000000)  * (10.000000,9.000000) = (-32.000000,116.000000)

Затем вы можете сравнить время двух разных стратегий умножения и выбрать, что лучше с вашим оборудованием.

Источник

Вы можете использовать cublasZdgmm.

cublasStatus_t cublasZdgmm(cublasHandle_t handle, cublasSideMode_t mode,
                      int m, int n,
                      const cuDoubleComplex *A, int lda,
                      const cuDoubleComplex *x, int incx,
                      cuDoubleComplex *C, int ldc)
Источник
Другие вопросы по тегам