sunwen/Aurora
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Fix complex fft bug

Browse Source
This commit is contained in:
kradchen
2023-12-14 17:42:06 +08:00
parent deaac03704
commit 5f906f78b8
2 changed files with 32 additions and 9 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
src/Function2D.cu
View File

@@ -1335,9 +1335,14 @@ CudaMatrix Aurora::fft(const CudaMatrix &aMatrix, long aFFTSize){
size_t needCopySize = (ColEleCount<aMatrix.getDimSize(0))?ColEleCount:aMatrix.getDimSize(0); size_t needCopySize = (ColEleCount<aMatrix.getDimSize(0))?ColEleCount:aMatrix.getDimSize(0);
size_t bufferSize = ColEleCount*aMatrix.getDimSize(1); size_t bufferSize = ColEleCount*aMatrix.getDimSize(1);
float* data = nullptr; float* data = nullptr;
cudaMalloc((void**)&data, sizeof(float)*2*bufferSize);
cudaMalloc((void**)&data, sizeof(float)*2*bufferSize);
if (aMatrix.isComplex()){
complexCopyKernel<<<aMatrix.getDimSize(1), 256>>>(aMatrix.getData(), data, needCopySize, aMatrix.getDimSize(0),ColEleCount);
}
else{
complexFillKernel<<<aMatrix.getDimSize(1), 256>>>(aMatrix.getData(), data, needCopySize, aMatrix.getDimSize(0),ColEleCount); complexFillKernel<<<aMatrix.getDimSize(1), 256>>>(aMatrix.getData(), data, needCopySize, aMatrix.getDimSize(0),ColEleCount);
}
auto ret = Aurora::CudaMatrix::fromRawData(data,ColEleCount,aMatrix.getDimSize(1),1,Complex); auto ret = Aurora::CudaMatrix::fromRawData(data,ColEleCount,aMatrix.getDimSize(1),1,Complex);
ExecFFT(ret,0); ExecFFT(ret,0);
return ret; return ret;

32
test/Function2D_Cuda_Test.cpp
View File

@@ -640,12 +640,6 @@ TEST_F(Function2D_Cuda_Test, fft)
EXPECT_FLOAT_EQ(fourDecimalRound(ifftResult[3].real()),2.0); EXPECT_FLOAT_EQ(fourDecimalRound(ifftResult[3].real()),2.0);
EXPECT_FLOAT_EQ(fourDecimalRound(ifftResult[11].real()),1.0); EXPECT_FLOAT_EQ(fourDecimalRound(ifftResult[11].real()),1.0);
EXPECT_FLOAT_EQ(fourDecimalRound(ifftResult[13].real()),2.0); EXPECT_FLOAT_EQ(fourDecimalRound(ifftResult[13].real()),2.0);
// Aurora::fftshift(fftrett);
// EXPECT_FLOAT_EQ(0.0729, fourDecimalRound(result[6].real()));
// EXPECT_FLOAT_EQ(2.4899, fourDecimalRound(result[7].imag()));
// EXPECT_FLOAT_EQ(-2., fourDecimalRound(result[10].real()));
// EXPECT_FLOAT_EQ(-0.2245, fourDecimalRound(result[11].imag()));
auto fftrett2 = Aurora::fft(ma,5); auto fftrett2 = Aurora::fft(ma,5);
auto fftrett2H =fftrett2.toHostMatrix(); auto fftrett2H =fftrett2.toHostMatrix();
@@ -668,7 +662,31 @@ TEST_F(Function2D_Cuda_Test, fft)
} }
{ {
float *input = new float[20]{1,1,0,2,2,0,1,1,0,2,1,1,0,2,2,0,1,1,0,2}; float *input = new float[20]{1,1,0,2,2,0,1,1,0,2,1,1,0,2,2,0,1,1,0,2};
auto ma = Aurora::Matrix::fromRawData(input,10,2); auto ma = Aurora::Matrix::fromRawData(input,10,1, 1,Aurora::Complex);
auto maD = ma.toDeviceMatrix();
auto ret1 = Aurora::fft(ma);
auto ret2 = Aurora::fft(maD);
for (size_t i = 0; i < ret1.getDataSize(); i++)
{
EXPECT_FLOAT_AE(ret1[i], ret2.getValue(i));
}
ret1 = Aurora::fft(ma,12);
ret2 = Aurora::fft(maD,12);
for (size_t i = 0; i < ret1.getDataSize(); i++)
{
EXPECT_FLOAT_AE(ret1[i], ret2.getValue(i));
}
ret1 = Aurora::fft(ma,8);
ret2 = Aurora::fft(maD,8);
for (size_t i = 0; i < ret1.getDataSize(); i++)
{
EXPECT_FLOAT_AE(ret1[i], ret2.getValue(i));
}
}
{
float *input = new float[20]{1,1,0,2,2,0,1,1,0,2,1,1,0,2,2,0,1,1,0,2};
auto ma = Aurora::Matrix::fromRawData(input,10,1,1, Aurora::Complex);
auto maD = ma.toDeviceMatrix(); auto maD = ma.toDeviceMatrix();
auto ret1 = Aurora::ifft(Aurora::fft(ma),12); auto ret1 = Aurora::ifft(Aurora::fft(ma),12);
auto ret2 = Aurora::ifft(Aurora::fft(maD),12); auto ret2 = Aurora::ifft(Aurora::fft(maD),12);