Commit f26f3901 authored by Pierre Paleo's avatar Pierre Paleo
Browse files

Remove dedicated halftomo kernel for cor > nx

parent 626f7541
......@@ -49,49 +49,3 @@ __global__ void halftomo_kernel(
}
}
// Same as previous kernel, but now the CoR is outside the image support
// i.e rotation_axis_position >= n_x
// Weigting is different !
__global__ void halftomo_kernel_cor_outside(
float* sinogram,
float* output,
float* weights,
int n_angles,
int n_x,
int rotation_axis_position
) {
int x = blockDim.x * blockIdx.x + threadIdx.x;
int y = blockDim.y * blockIdx.y + threadIdx.y;
int n_a2 = n_angles / 2;
int d = rotation_axis_position - n_x;
int n_x2 = 2 * rotation_axis_position;
if ((x >= n_x2) || (y >= n_a2)) return;
// output[:, :nx] = sino[:n_a2, :nx]
if (x < n_x) {
output[y * n_x2 + x] = sinogram[y * n_x + x];
}
// output[:, nx : nx + d] = (1 - weights) * sino[:n_a2, -d :][:, ::-1]
else if (x < n_x + d) {
float w = weights[x - n_x];
// output[y * n_x2 + x] = (1.0f - w) * sinogram[y*n_x + n_x - 1 - (n_x - x)];
output[y * n_x2 + x] = (1.0f - w) * sinogram[y*n_x + 2 * n_x - 1 - x];
}
// output[:, nx + d : nx + 2*d] = weights * sino[n_a2:, ::-1][:, :d] = sino[n_a2:, -d-1:-n_x-1:-1]
else if (x < n_x + 2 * d) {
float w = weights[x - (n_x + d)];
output[y * n_x2 + x] = w * sinogram[(n_a2 + y)*n_x + x - 2*d];
}
// output[:, nx+2*d:] = sino[n_a2:, ::-1][:, 2 * d :] = sino[n_a2:, -2*d-1:-n_x-1:-1]
else {
output[y * n_x2 + x] = sinogram[(n_a2 + y)*n_x + (n_x2 - 1 - x)];
}
}
\ No newline at end of file
......@@ -27,10 +27,7 @@ class CudaSinoProcessing(SinoProcessing, CudaProcessing):
def _init_cuda_halftomo(self):
if not(self.halftomo):
return
if self._rot_center_int < self.n_x:
kernel_name = "halftomo_kernel"
else:
kernel_name = "halftomo_kernel_cor_outside"
kernel_name = "halftomo_kernel"
self.halftomo_kernel = CudaKernel(
kernel_name,
get_cuda_srcfile("halftomo.cu"),
......@@ -93,7 +90,7 @@ class CudaSinoProcessing(SinoProcessing, CudaProcessing):
else:
sinos = garray.zeros(self.sinos_halftomo_shape, dtype=np.float32)
# FIXME: TEMPORARY PATCH
# FIXME: TEMPORARY PATCH, waiting for cuda implementation
if self._rot_center_int > self.n_x:
return self._radios_to_sinos_halftomo_external_cor(radios, sinos)
#
......@@ -125,7 +122,7 @@ class CudaSinoProcessing(SinoProcessing, CudaProcessing):
def _radios_to_sinos_halftomo_external_cor(self, radios, sinos):
"""
TEMPORARY PATCH waiting to fix cuda kernel
TEMPORARY PATCH waiting to have a cuda implementation
Processing is done by getting reach radio on host, which is suboptimal
"""
n_a, n_z, n_x = radios.shape
......
Supports Markdown
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment