summaryrefslogtreecommitdiffstats
diff options
context:
space:
mode:
authorWillem Jan Palenstijn <wjp@usecode.org>2021-01-22 11:48:32 +0100
committerWillem Jan Palenstijn <wjp@usecode.org>2021-01-22 11:48:32 +0100
commitc20e4b3fe683fa27ab3dcba03a3f5cdfb9b8db0f (patch)
tree4eb30596204a4571417270b969e9a7ae80b6d6e6
parent74f822299c0292a77a0dfdf79a8cee4a2d4b6c1d (diff)
downloadastra-c20e4b3fe683fa27ab3dcba03a3f5cdfb9b8db0f.tar.gz
astra-c20e4b3fe683fa27ab3dcba03a3f5cdfb9b8db0f.tar.bz2
astra-c20e4b3fe683fa27ab3dcba03a3f5cdfb9b8db0f.tar.xz
astra-c20e4b3fe683fa27ab3dcba03a3f5cdfb9b8db0f.zip
Fix supersampling version of cone_bp and add test
-rw-r--r--cuda/3d/cone_bp.cu8
-rw-r--r--tests/python/test_rec_scaling.py27
2 files changed, 28 insertions, 7 deletions
diff --git a/cuda/3d/cone_bp.cu b/cuda/3d/cone_bp.cu
index df1e19f..3525eb4 100644
--- a/cuda/3d/cone_bp.cu
+++ b/cuda/3d/cone_bp.cu
@@ -213,15 +213,15 @@ __global__ void dev_cone_BP_SS(void* D_volData, unsigned int volPitch, int start
float fZs = fZ;
for (int iSubZ = 0; iSubZ < iRaysPerVoxelDim; ++iSubZ) {
- const float fUNum = fCu.w + fX * fCu.x + fY * fCu.y + fZ * fCu.z;
- const float fVNum = fCv.w + fX * fCv.x + fY * fCv.y + fZ * fCv.z;
- const float fDen = fCd.w + fX * fCd.x + fY * fCd.y + fZ * fCd.z;
+ const float fUNum = fCu.w + fXs * fCu.x + fYs * fCu.y + fZs * fCu.z;
+ const float fVNum = fCv.w + fXs * fCv.x + fYs * fCv.y + fZs * fCv.z;
+ const float fDen = fCd.w + fXs * fCd.x + fYs * fCd.y + fZs * fCd.z;
const float fr = __fdividef(1.0f, fDen);
const float fU = fUNum * fr;
const float fV = fVNum * fr;
- fVal += tex3D(gT_coneProjTexture, fU, fV, fAngle) * fr;
+ fVal += tex3D(gT_coneProjTexture, fU, fAngle, fV) * fr * fr;
fZs += fSubStep;
}
diff --git a/tests/python/test_rec_scaling.py b/tests/python/test_rec_scaling.py
index 621fd8a..ee27efc 100644
--- a/tests/python/test_rec_scaling.py
+++ b/tests/python/test_rec_scaling.py
@@ -76,7 +76,7 @@ def ProjectionGeometries(type):
class TestRecScale(unittest.TestCase):
- def single_test(self, geom_type, proj_type, alg, iters):
+ def single_test(self, geom_type, proj_type, alg, iters, vss, dss):
if alg == 'FBP' and 'fanflat' in geom_type:
self.skipTest('CPU FBP is parallel-beam only')
is3D = (geom_type in ['parallel3d', 'cone'])
@@ -88,7 +88,12 @@ class TestRecScale(unittest.TestCase):
else:
vol = np.zeros((64,64,64),dtype=np.float32)
vol[25:35,25:35,25:35] = 1
- proj_id = astra.create_projector(proj_type, pg, vg)
+ options = {}
+ if vss > 1:
+ options["VoxelSuperSampling"] = vss
+ if dss > 1:
+ options["DetectorSuperSampling"] = vss
+ proj_id = astra.create_projector(proj_type, pg, vg, options=options)
if not is3D:
sino_id, sinogram = astra.create_sino(vol, proj_id)
else:
@@ -185,6 +190,14 @@ __algs_cone = {
}
+__combinations_ss = {
+ 'parallel': [ { 'projector': 'cuda', 'alg': 'SIRT_CUDA', 'iters': 50 } ],
+ 'fanflat': [ { 'projector': 'cuda', 'alg': 'SIRT_CUDA', 'iters': 50 } ],
+ 'parallel3d': [ { 'projector': 'cuda3d', 'alg': 'SIRT3D_CUDA', 'iters': 200 } ],
+ 'cone': [ { 'projector': 'cuda3d', 'alg': 'SIRT3D_CUDA', 'iters': 200 } ]
+}
+
+
for k, l in __combinations.items():
for v in l:
@@ -199,7 +212,7 @@ for k, l in __combinations.items():
A = __algs
for a, i in A.items():
def f(k, v, a, i):
- return lambda self: self.single_test(k, v, a, i)
+ return lambda self: self.single_test(k, v, a, i, 1, 1)
setattr(TestRecScale, 'test_' + a + '_' + k + '_' + v, f(k,v,a,i))
for k, l in __combinations_adjoint.items():
@@ -208,6 +221,14 @@ for k, l in __combinations_adjoint.items():
return lambda self: self.single_test_adjoint3D(k, v)
setattr(TestRecScale, 'test_adjoint_' + k + '_' + v, g(k,v))
+for k, l in __combinations_ss.items():
+ for A in l:
+ for vss in [1, 2]:
+ for dss in [1, 2]:
+ def h(k, v, a, i, vss, dss):
+ return lambda self: self.single_test(k, v, a, i, vss, dss)
+ setattr(TestRecScale, 'test_ss_' + a + '_' + k + '_' + v + '_' + str(vss) + '_' + str(dss), h(k, A['projector'], A['alg'], A['iters'], vss, dss))
+
if __name__ == '__main__':
unittest.main()