TEST(Projection, Forward)
{
cudaDeviceReset();
//Case 1:
{
int3 dimsvolume = {8, 8, 8};
int3 dimsimage = {8, 8, 1};
tfloat3 angles = tfloat3(PI / 2.0f, PI / 2.0f, 0.0f);
tfloat2 shifts = tfloat2(0, 0);
tfloat weight = (tfloat)1;
tfloat* d_inputvolume = (tfloat*)CudaMallocFromBinaryFile("Data\\Projection\\Input_Forward_1.bin");
tfloat* d_volumepsf = CudaMallocValueFilled(ElementsFFT(dimsvolume), (tfloat)1);
tfloat* d_proj = (tfloat*)CudaMallocValueFilled(Elements(dimsimage), (tfloat)0);
tfloat* desired_output = (tfloat*)MallocFromBinaryFile("Data\\Projection\\Output_Forward_1.bin");
tfloat* d_projpsf = CudaMallocValueFilled(ElementsFFT(dimsimage), (tfloat)0);
d_ProjForward(d_inputvolume, d_volumepsf, dimsvolume, d_proj, d_projpsf, &angles, &shifts, T_INTERP_CUBIC, 1);
tfloat* h_output = (tfloat*)MallocFromDeviceArray(d_proj, Elements(dimsimage) * sizeof(tfloat));
double MeanAbsolute = GetMeanAbsoluteError((tfloat*)desired_output, (tfloat*)h_output, Elements(dimsimage));
ASSERT_LE(MeanAbsolute, 1e-5);
cudaFree(d_inputvolume);
cudaFree(d_proj);
free(desired_output);
free(h_output);
}
cudaDeviceReset();
}
TEST(Transformation, Bin)
{
for(int i = 11; i < 12; i++)
{
cudaDeviceReset();
srand(i);
int size = (1<<i);
int batch = 1;
int bincount = 5;
tfloat* h_input = (tfloat*)malloc(size * size * batch * sizeof(tfloat));
for(int b = 0; b < batch; b++)
{
for(int j = 0; j < size * size; j++)
h_input[b * size * size + j] = (tfloat)(j % (1<<bincount));
}
tfloat* d_input = (tfloat*)CudaMallocFromHostArray(h_input, size * size * batch * sizeof(tfloat));
tfloat* d_result;
cudaMalloc((void**)&d_result, size * size / (1<<(bincount * 2)) * sizeof(tfloat));
int3 dims;
dims.x = size;
dims.y = size;
d_Bin(d_input, d_result, dims, bincount, 1);
tfloat* h_result = (tfloat*)MallocFromDeviceArray(d_result, size * size / (1<<(bincount * 2)) * sizeof(tfloat));
ASSERT_ARRAY_EQ(h_result, (tfloat)((1<<bincount) - 1) / (tfloat)2, size * size / (1<<(bincount * 2)));
cudaFree(d_input);
cudaFree(d_result);
free(h_input);
free(h_result);
cudaDeviceReset();
}
for(int i = 9; i < 10; i++)
{
cudaDeviceReset();
srand(i);
size_t size = (1<<i);
size_t batch = 1;
size_t bincount = 2;
tfloat* h_input;
cudaMallocHost((void**)&h_input, size * size * size * batch * sizeof(tfloat), 0);
for(int b = 0; b < batch; b++)
{
for(int j = 0; j < size * size * size; j++)
h_input[b * size * size * size + j] = (tfloat)(j % (1<<bincount));
}
tfloat* d_input = (tfloat*)CudaMallocFromHostArray(h_input, size * size * size * batch * sizeof(tfloat));
tfloat* d_result;
cudaMalloc((void**)&d_result, size * size * size / (1<<(bincount * 3)) * batch * sizeof(tfloat));
int3 dims;
dims.x = size;
dims.y = size;
dims.z = size;
d_Bin(d_input, d_result, dims, bincount, batch);
tfloat* h_result = (tfloat*)MallocFromDeviceArray(d_result, size * size * size / (1<<(bincount * 3)) * batch * sizeof(tfloat));
ASSERT_ARRAY_EQ(h_result, (tfloat)((1<<bincount) - 1) / (tfloat)2, size * size / (1<<(bincount * 2)));
cudaFreeHost(h_input);
free(h_result);
cudaFree(d_input);
cudaFree(d_result);
cudaDeviceReset();
}
}
TEST(Correlation, Peak)
{
cudaDeviceReset();
//Case 1:
{
int3 dims = {8, 8, 1};
tfloat* d_input = (tfloat*)CudaMallocFromBinaryFile("Data\\Correlation\\Input_Peak_1.bin");
tfloat3* desired_output = (tfloat3*)MallocFromBinaryFile("Data\\Correlation\\Output_Peak_1.bin");
tfloat3* d_positions;
cudaMalloc((void**)&d_positions, sizeof(tfloat3));
tfloat* d_values;
cudaMalloc((void**)&d_values, sizeof(tfloat));
d_Peak(d_input, d_positions, d_values, dims, T_PEAK_MODE::T_PEAK_INTEGER);
tfloat3* h_positions = (tfloat3*)MallocFromDeviceArray(d_positions, sizeof(tfloat3));
tfloat* h_values = (tfloat*)MallocFromDeviceArray(d_values, sizeof(tfloat));
double MeanRelative = GetMeanRelativeError((tfloat*)desired_output, (tfloat*)h_positions, 3);
ASSERT_LE(MeanRelative, 1e-5);
cudaFree(d_input);
cudaFree(d_positions);
cudaFree(d_values);
free(desired_output);
free(h_positions);
free(h_values);
}
//Case 2:
{
int3 dims = {20, 20, 1};
tfloat* d_input = (tfloat*)CudaMallocFromBinaryFile("Data\\Correlation\\Input_Peak_2.bin");
tfloat3* desired_output = (tfloat3*)MallocFromBinaryFile("Data\\Correlation\\Output_Peak_2.bin");
tfloat3* d_positions;
cudaMalloc((void**)&d_positions, sizeof(tfloat3));
tfloat* d_values;
cudaMalloc((void**)&d_values, sizeof(tfloat));
d_Peak(d_input, d_positions, d_values, dims, T_PEAK_MODE::T_PEAK_SUBCOARSE);
tfloat3* h_positions = (tfloat3*)MallocFromDeviceArray(d_positions, sizeof(tfloat3));
tfloat* h_values = (tfloat*)MallocFromDeviceArray(d_values, sizeof(tfloat));
double MeanRelative = GetMeanAbsoluteError((tfloat*)desired_output, (tfloat*)h_positions, DimensionCount(dims));
ASSERT_LE(MeanRelative, 0.05);
cudaFree(d_input);
cudaFree(d_positions);
cudaFree(d_values);
free(desired_output);
free(h_positions);
free(h_values);
}
//Case 3:
{
int3 dims = {20, 20, 1};
tfloat* d_input = (tfloat*)CudaMallocFromBinaryFile("Data\\Correlation\\Input_Peak_2.bin");
tfloat3* desired_output = (tfloat3*)MallocFromBinaryFile("Data\\Correlation\\Output_Peak_2.bin");
tfloat3* d_positions;
cudaMalloc((void**)&d_positions, sizeof(tfloat3));
tfloat* d_values;
cudaMalloc((void**)&d_values, sizeof(tfloat));
d_Peak(d_input, d_positions, d_values, dims, T_PEAK_MODE::T_PEAK_SUBFINE);
tfloat3* h_positions = (tfloat3*)MallocFromDeviceArray(d_positions, sizeof(tfloat3));
tfloat* h_values = (tfloat*)MallocFromDeviceArray(d_values, sizeof(tfloat));
double MeanRelative = GetMeanAbsoluteError((tfloat*)desired_output, (tfloat*)h_positions, DimensionCount(dims));
ASSERT_LE(MeanRelative, 0.02);
cudaFree(d_input);
cudaFree(d_positions);
cudaFree(d_values);
free(desired_output);
free(h_positions);
free(h_values);
}
cudaDeviceReset();
}
