summaryrefslogtreecommitdiffstats
path: root/cuda
diff options
context:
space:
mode:
authorWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2021-12-02 11:52:10 +0100
committerWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2021-12-06 11:01:36 +0100
commit33f52988134d11096b69352671c54045a30a82d4 (patch)
tree92d810de0471deb5f78ac7fa87213c9f33caafb2 /cuda
parent86615d4161b050fbf3335e30ae85801aa1cefe92 (diff)
downloadastra-33f52988134d11096b69352671c54045a30a82d4.tar.gz
astra-33f52988134d11096b69352671c54045a30a82d4.tar.bz2
astra-33f52988134d11096b69352671c54045a30a82d4.tar.xz
astra-33f52988134d11096b69352671c54045a30a82d4.zip
Add transferConstants functions for FP kernels for consistency
Diffstat (limited to 'cuda')
-rw-r--r--cuda/3d/cone_fp.cu51
-rw-r--r--cuda/3d/par3d_fp.cu49
2 files changed, 58 insertions, 42 deletions
diff --git a/cuda/3d/cone_fp.cu b/cuda/3d/cone_fp.cu
index 2ef58ee..ccdaf16 100644
--- a/cuda/3d/cone_fp.cu
+++ b/cuda/3d/cone_fp.cu
@@ -114,6 +114,34 @@ struct SCALE_NONCUBE {
};
+bool transferConstants(const SConeProjection* angles, unsigned int iProjAngles)
+{
+ // transfer angles to constant memory
+ float* tmp = new float[iProjAngles];
+
+#define TRANSFER_TO_CONSTANT(name) do { for (unsigned int i = 0; i < iProjAngles; ++i) tmp[i] = angles[i].f##name ; cudaMemcpyToSymbol(gC_##name, tmp, iProjAngles*sizeof(float), 0, cudaMemcpyHostToDevice); } while (0)
+
+ TRANSFER_TO_CONSTANT(SrcX);
+ TRANSFER_TO_CONSTANT(SrcY);
+ TRANSFER_TO_CONSTANT(SrcZ);
+ TRANSFER_TO_CONSTANT(DetSX);
+ TRANSFER_TO_CONSTANT(DetSY);
+ TRANSFER_TO_CONSTANT(DetSZ);
+ TRANSFER_TO_CONSTANT(DetUX);
+ TRANSFER_TO_CONSTANT(DetUY);
+ TRANSFER_TO_CONSTANT(DetUZ);
+ TRANSFER_TO_CONSTANT(DetVX);
+ TRANSFER_TO_CONSTANT(DetVY);
+ TRANSFER_TO_CONSTANT(DetVZ);
+
+#undef TRANSFER_TO_CONSTANT
+
+ delete[] tmp;
+
+ return true;
+}
+
+
// threadIdx: x = ??? detector (u?)
// y = relative angle
@@ -298,27 +326,8 @@ bool ConeFP_Array_internal(cudaPitchedPtr D_projData,
unsigned int angleCount, const SConeProjection* angles,
const SProjectorParams3D& params)
{
- // transfer angles to constant memory
- float* tmp = new float[angleCount];
-
-#define TRANSFER_TO_CONSTANT(name) do { for (unsigned int i = 0; i < angleCount; ++i) tmp[i] = angles[i].f##name ; cudaMemcpyToSymbol(gC_##name, tmp, angleCount*sizeof(float), 0, cudaMemcpyHostToDevice); } while (0)
-
- TRANSFER_TO_CONSTANT(SrcX);
- TRANSFER_TO_CONSTANT(SrcY);
- TRANSFER_TO_CONSTANT(SrcZ);
- TRANSFER_TO_CONSTANT(DetSX);
- TRANSFER_TO_CONSTANT(DetSY);
- TRANSFER_TO_CONSTANT(DetSZ);
- TRANSFER_TO_CONSTANT(DetUX);
- TRANSFER_TO_CONSTANT(DetUY);
- TRANSFER_TO_CONSTANT(DetUZ);
- TRANSFER_TO_CONSTANT(DetVX);
- TRANSFER_TO_CONSTANT(DetVY);
- TRANSFER_TO_CONSTANT(DetVZ);
-
-#undef TRANSFER_TO_CONSTANT
-
- delete[] tmp;
+ if (!transferConstants(angles, angleCount))
+ return false;
std::list<cudaStream_t> streams;
dim3 dimBlock(g_detBlockU, g_anglesPerBlock); // region size, angles
diff --git a/cuda/3d/par3d_fp.cu b/cuda/3d/par3d_fp.cu
index 9475897..fda6f93 100644
--- a/cuda/3d/par3d_fp.cu
+++ b/cuda/3d/par3d_fp.cu
@@ -113,6 +113,32 @@ struct SCALE_NONCUBE {
__device__ float scale(float a1, float a2) const { return sqrt(a1*a1*fScale1+a2*a2*fScale2+1.0f) * fOutputScale; }
};
+bool transferConstants(const SPar3DProjection* angles, unsigned int iProjAngles)
+{
+ // transfer angles to constant memory
+ float* tmp = new float[iProjAngles];
+
+#define TRANSFER_TO_CONSTANT(name) do { for (unsigned int i = 0; i < iProjAngles; ++i) tmp[i] = angles[i].f##name ; cudaMemcpyToSymbol(gC_##name, tmp, iProjAngles*sizeof(float), 0, cudaMemcpyHostToDevice); } while (0)
+
+ TRANSFER_TO_CONSTANT(RayX);
+ TRANSFER_TO_CONSTANT(RayY);
+ TRANSFER_TO_CONSTANT(RayZ);
+ TRANSFER_TO_CONSTANT(DetSX);
+ TRANSFER_TO_CONSTANT(DetSY);
+ TRANSFER_TO_CONSTANT(DetSZ);
+ TRANSFER_TO_CONSTANT(DetUX);
+ TRANSFER_TO_CONSTANT(DetUY);
+ TRANSFER_TO_CONSTANT(DetUZ);
+ TRANSFER_TO_CONSTANT(DetVX);
+ TRANSFER_TO_CONSTANT(DetVY);
+ TRANSFER_TO_CONSTANT(DetVZ);
+
+#undef TRANSFER_TO_CONSTANT
+
+ delete[] tmp;
+
+ return true;
+}
// threadIdx: x = u detector
@@ -400,27 +426,8 @@ bool Par3DFP_Array_internal(cudaPitchedPtr D_projData,
unsigned int angleCount, const SPar3DProjection* angles,
const SProjectorParams3D& params)
{
- // transfer angles to constant memory
- float* tmp = new float[angleCount];
-
-#define TRANSFER_TO_CONSTANT(name) do { for (unsigned int i = 0; i < angleCount; ++i) tmp[i] = angles[i].f##name ; cudaMemcpyToSymbol(gC_##name, tmp, angleCount*sizeof(float), 0, cudaMemcpyHostToDevice); } while (0)
-
- TRANSFER_TO_CONSTANT(RayX);
- TRANSFER_TO_CONSTANT(RayY);
- TRANSFER_TO_CONSTANT(RayZ);
- TRANSFER_TO_CONSTANT(DetSX);
- TRANSFER_TO_CONSTANT(DetSY);
- TRANSFER_TO_CONSTANT(DetSZ);
- TRANSFER_TO_CONSTANT(DetUX);
- TRANSFER_TO_CONSTANT(DetUY);
- TRANSFER_TO_CONSTANT(DetUZ);
- TRANSFER_TO_CONSTANT(DetVX);
- TRANSFER_TO_CONSTANT(DetVY);
- TRANSFER_TO_CONSTANT(DetVZ);
-
-#undef TRANSFER_TO_CONSTANT
-
- delete[] tmp;
+ if (!transferConstants(angles, angleCount))
+ return false;
std::list<cudaStream_t> streams;
dim3 dimBlock(g_detBlockU, g_anglesPerBlock); // region size, angles