summaryrefslogtreecommitdiffstats
path: root/cuda/2d/arith.cu
diff options
context:
space:
mode:
authorWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2014-10-02 13:50:59 +0200
committerWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2014-10-02 13:51:54 +0200
commit12fd6925502d3ead9b73498f44f9b2f9abbc9bea (patch)
tree10970caff57e8f6b127dc223f73270eb3a490dc8 /cuda/2d/arith.cu
parent8188b0697751d7c7ec6fe3069d0fceb50ebd0c9e (diff)
downloadastra-12fd6925502d3ead9b73498f44f9b2f9abbc9bea.tar.gz
astra-12fd6925502d3ead9b73498f44f9b2f9abbc9bea.tar.bz2
astra-12fd6925502d3ead9b73498f44f9b2f9abbc9bea.tar.xz
astra-12fd6925502d3ead9b73498f44f9b2f9abbc9bea.zip
Add CUDA SIRT::doSlabCorrections() function
This function optionally compensates for effectively infinitely large slab-like objects of finite thickness 1.
Diffstat (limited to 'cuda/2d/arith.cu')
-rw-r--r--cuda/2d/arith.cu9
1 files changed, 9 insertions, 0 deletions
diff --git a/cuda/2d/arith.cu b/cuda/2d/arith.cu
index 04d4de9..93b1ee0 100644
--- a/cuda/2d/arith.cu
+++ b/cuda/2d/arith.cu
@@ -68,6 +68,14 @@ struct opMul {
out *= in;
}
};
+struct opDiv {
+ __device__ void operator()(float& out, const float in) {
+ if (in > 0.000001f) // out is assumed to be positive
+ out /= in;
+ else
+ out = 0.0f;
+ }
+};
struct opMul2 {
__device__ void operator()(float& out, const float in1, const float in2) {
out *= in1 * in2;
@@ -682,6 +690,7 @@ INST_DtoD(opDividedBy)
INST_toD(opInvert)
INST_FtoD(opSet)
INST_FtoD(opMul)
+INST_DtoD(opDiv)
INST_DFtoD(opMulMask)
INST_FtoD(opAdd)
INST_FtoD(opClampMin)