summaryrefslogtreecommitdiffstats
path: root/cuda/3d
diff options
context:
space:
mode:
Diffstat (limited to 'cuda/3d')
-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