TEST(Reconstruction, Fourier) { cudaDeviceReset(); //Case 1: /*{ int3 dimsvolume = {16, 16, 16}; int3 dimsimage = {16, 16, 2249}; tfloat* d_inputproj = (tfloat*)CudaMallocFromBinaryFile("Data\\Reconstruction\\Input_ARTProj_2.bin"); tfloat3* h_inputangles = (tfloat3*)MallocFromBinaryFile("Data\\Reconstruction\\Input_ARTAngles_2.bin"); tfloat* desired_output = (tfloat*)MallocFromBinaryFile("Data\\Reconstruction\\Output_ART_2.bin"); tfloat* d_volume; cudaMalloc((void**)&d_volume, Elements(dimsvolume) * sizeof(tfloat)); tfloat* h_output = (tfloat*)MallocFromDeviceArray(d_volume, Elements(dimsvolume) * sizeof(tfloat)); tfloat outputmax = (tfloat)-999999; for (int i = 0; i < Elements(dimsvolume); i++) outputmax = max(outputmax, h_output[i]); //for (int i = 0; i < Elements(dimsvolume); i++) //h_output[i] /= outputmax; double MeanAbsolute = GetMeanAbsoluteError((tfloat*)desired_output, (tfloat*)h_output, Elements(dimsvolume)); //ASSERT_LE(MeanRelative, 1e-5); cudaFree(d_volume); cudaFree(d_inputproj); free(desired_output); free(h_output); }*/ //Case 2: { int3 dimsori = toInt3(128, 128, 128); int3 dimspadded = toInt3(259, 259, 259); tfloat* h_weights = (tfloat*)malloc(ElementsFFT(dimspadded) * sizeof(tfloat)); ReadMRC("d_weights.mrc", (void**)&h_weights); tfloat* d_weights = (tfloat*)CudaMallocFromHostArray(h_weights, ElementsFFT(dimspadded) * sizeof(tfloat)); tcomplex* h_data = (tcomplex*)malloc(ElementsFFT(dimspadded) * sizeof(tcomplex)); ReadMRC("d_dataRe.mrc", (void**)&h_weights); for (int i = 0; i < ElementsFFT(dimspadded); i++) h_data[i].x = h_weights[i]; ReadMRC("d_dataIm.mrc", (void**)&h_weights); for (int i = 0; i < ElementsFFT(dimspadded); i++) h_data[i].y = h_weights[i]; tcomplex* d_data = (tcomplex*)CudaMallocFromHostArray(h_data, ElementsFFT(dimspadded) * sizeof(tcomplex)); tfloat* d_reconstructed = CudaMallocValueFilled(Elements(dimsori), (tfloat)0); d_ReconstructGridding(d_data, d_weights, d_reconstructed, dimsori, dimspadded); } cudaDeviceReset(); }
TEST(Transformation, Warp2D) { cudaDeviceReset(); //Case 1: { int2 dimsimage = toInt2(1024, 1024); int nframes = 40; tfloat* d_input = (tfloat*)CudaMallocFromBinaryFile("Data\\Transformation\\Input_Warp2D.bin"); tfloat* d_output = CudaMallocValueFilled(Elements2(dimsimage) * nframes, (tfloat)0); int2 dimsgrid = toInt2(4, 4); tfloat2* h_grid = (tfloat2*)MallocValueFilled(Elements2(dimsgrid) * 2, (tfloat)0); for (uint n = 0; n < nframes; n++) { h_grid[1 * 4 + 0] = tfloat2(n * 1.0, 0.0); h_grid[1 * 4 + 1] = tfloat2(n * 1.0, 0.0); h_grid[1 * 4 + 2] = tfloat2(n * 1.0, 0.0); h_grid[1 * 4 + 3] = tfloat2(n * 1.0, 0.0); h_grid[2 * 4 + 0] = tfloat2(n * 1.0, 0.0); h_grid[2 * 4 + 1] = tfloat2(n * 1.0, 0.0); h_grid[2 * 4 + 2] = tfloat2(n * 1.0, 0.0); h_grid[2 * 4 + 3] = tfloat2(n * 1.0, 0.0); tfloat2* d_grid = (tfloat2*)CudaMallocFromHostArray(h_grid, Elements2(dimsgrid) * sizeof(tfloat2)); d_Warp2D(d_input, dimsimage, d_grid, dimsgrid, d_output + Elements2(dimsimage) * n); cudaFree(d_grid); } d_WriteMRC(d_output, toInt3(dimsimage.x, dimsimage.y, nframes), "d_warped.mrc"); } 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(); } }