|
@@ -46,3 +46,36 @@
|
|
|
id.y * get_global_size (0) + id.x] = read_imagef (projection, sampler, pixel).x;
|
|
|
}
|
|
|
}
|
|
|
+
|
|
|
+
|
|
|
+__kernel void backproject_from_3d (global float *volume,
|
|
|
+ __read_only image3d_t projections,
|
|
|
+ const sampler_t sampler,
|
|
|
+ const float2 x_region,
|
|
|
+ const float2 y_region,
|
|
|
+ const float2 z_region,
|
|
|
+ constant float *sines,
|
|
|
+ constant float *cosines,
|
|
|
+ const float2 center,
|
|
|
+ const float sin_lamino,
|
|
|
+ const float cos_lamino)
|
|
|
+{
|
|
|
+ int3 id = (int3) (get_global_id (0), get_global_id (1), get_global_id (2));
|
|
|
+ int i;
|
|
|
+ float result = 0.0f;
|
|
|
+ float4 pixel;
|
|
|
+ float3 voxel;
|
|
|
+
|
|
|
+ voxel.x = mad((float) id.x, x_region.y, x_region.x);
|
|
|
+ voxel.y = mad((float) id.y, y_region.y, y_region.x);
|
|
|
+ voxel.z = mad((float) id.z, z_region.y, z_region.x);
|
|
|
+
|
|
|
+ for (i = 0; i < get_image_depth (projections); i++) {
|
|
|
+ pixel.x = voxel.x * cosines[i] + voxel.y * sines[i] + center.x;
|
|
|
+ pixel.y = voxel.x * cos_lamino * sines[i] - voxel.y * cos_lamino * cosines[i] + voxel.z * sin_lamino + center.y;
|
|
|
+ pixel.z = (float) i;
|
|
|
+ result += read_imagef (projections, sampler, pixel).x;
|
|
|
+ }
|
|
|
+
|
|
|
+ volume[id.z * get_global_size (0) * get_global_size (1) + id.y * get_global_size (0) + id.x] = result;
|
|
|
+}
|