我已经用CUDA构建了一个简单的内核,以执行两个复向量的逐元素向量乘法。内核代码如下(
作为最后一招,我尝试通过在共享内存中加载两个向量,然后从那里工作来优化这段代码(请参见下面的
所以我的问题是:
multiplyElementwise
)。它能正常工作,但是我注意到像缩放向量这样的其他看似简单的操作在像CUBLAS或CULA这样的库中进行了优化,所以我想知道是否可能用库调用替换我的代码?令我惊讶的是,无论是CUBLAS还是CULA都没有这个选项,我试图通过使其中一个向量成为对角线矩阵-向量乘积的对角线来模拟它,但结果非常慢。作为最后一招,我尝试通过在共享内存中加载两个向量,然后从那里工作来优化这段代码(请参见下面的
multiplyElementwiseFast
),但这比我的原始代码更慢。所以我的问题是:
- 是否有可以进行元素级向量-向量乘法的库?
- 如果没有,我能加速我的代码(
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];
}
}