summaryrefslogtreecommitdiffstats
path: root/cuda/3d/cone_fp.cu
diff options
context:
space:
mode:
authorWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2014-06-23 10:08:37 +0000
committerwpalenst <Willem.Jan.Palenstijn@cwi.nl>2014-06-23 10:08:37 +0000
commit00767fae7142a66b182508448951816f9c95f189 (patch)
tree9f55239cd8960ce703aabc513f765afdbb91374e /cuda/3d/cone_fp.cu
parent97ba7288f6d665c4442b3c9873128529c7dcf508 (diff)
downloadastra-00767fae7142a66b182508448951816f9c95f189.tar.gz
astra-00767fae7142a66b182508448951816f9c95f189.tar.bz2
astra-00767fae7142a66b182508448951816f9c95f189.tar.xz
astra-00767fae7142a66b182508448951816f9c95f189.zip
Remove angle limits in cone
Diffstat (limited to 'cuda/3d/cone_fp.cu')
-rw-r--r--cuda/3d/cone_fp.cu37
1 files changed, 25 insertions, 12 deletions
diff --git a/cuda/3d/cone_fp.cu b/cuda/3d/cone_fp.cu
index 0b1f012..d049151 100644
--- a/cuda/3d/cone_fp.cu
+++ b/cuda/3d/cone_fp.cu
@@ -299,18 +299,14 @@ __global__ void cone_FP_SS_t(float* D_projData, unsigned int projPitch,
}
-
-bool ConeFP_Array(cudaArray *D_volArray,
- cudaPitchedPtr D_projData,
- const SDimensions3D& dims, const SConeProjection* angles,
+bool ConeFP_Array_internal(cudaPitchedPtr D_projData,
+ const SDimensions3D& dims, unsigned int angleCount, const SConeProjection* angles,
float fOutputScale)
{
- bindVolumeDataTexture(D_volArray);
-
// transfer angles to constant memory
- float* tmp = new float[dims.iProjAngles];
+ float* tmp = new float[angleCount];
-#define TRANSFER_TO_CONSTANT(name) do { for (unsigned int i = 0; i < dims.iProjAngles; ++i) tmp[i] = angles[i].f##name ; cudaMemcpyToSymbol(gC_##name, tmp, dims.iProjAngles*sizeof(float), 0, cudaMemcpyHostToDevice); } while (0)
+#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);
@@ -343,9 +339,9 @@ bool ConeFP_Array(cudaArray *D_volArray,
// timeval t;
// tic(t);
- for (unsigned int a = 0; a <= dims.iProjAngles; ++a) {
+ for (unsigned int a = 0; a <= angleCount; ++a) {
int dir;
- if (a != dims.iProjAngles) {
+ if (a != angleCount) {
float dX = fabsf(angles[a].fSrcX - (angles[a].fDetSX + dims.iProjU*angles[a].fDetUX*0.5f + dims.iProjV*angles[a].fDetVX*0.5f));
float dY = fabsf(angles[a].fSrcY - (angles[a].fDetSY + dims.iProjU*angles[a].fDetUY*0.5f + dims.iProjV*angles[a].fDetVY*0.5f));
float dZ = fabsf(angles[a].fSrcZ - (angles[a].fDetSZ + dims.iProjU*angles[a].fDetUZ*0.5f + dims.iProjV*angles[a].fDetVZ*0.5f));
@@ -358,7 +354,7 @@ bool ConeFP_Array(cudaArray *D_volArray,
dir = 2;
}
- if (a == dims.iProjAngles || dir != blockDirection) {
+ if (a == angleCount || dir != blockDirection) {
// block done
blockEnd = a;
@@ -414,6 +410,7 @@ bool ConeFP_Array(cudaArray *D_volArray,
return true;
}
+
bool ConeFP(cudaPitchedPtr D_volumeData,
cudaPitchedPtr D_projData,
const SDimensions3D& dims, const SConeProjection* angles,
@@ -423,8 +420,24 @@ bool ConeFP(cudaPitchedPtr D_volumeData,
cudaArray* cuArray = allocateVolumeArray(dims);
transferVolumeToArray(D_volumeData, cuArray, dims);
+ bindVolumeDataTexture(cuArray);
+
+ bool ret;
- bool ret = ConeFP_Array(cuArray, D_projData, dims, angles, fOutputScale);
+ for (unsigned int iAngle = 0; iAngle < dims.iProjAngles; iAngle += g_MaxAngles) {
+ unsigned int iEndAngle = iAngle + g_MaxAngles;
+ if (iEndAngle >= dims.iProjAngles)
+ iEndAngle = dims.iProjAngles;
+
+ cudaPitchedPtr D_subprojData = D_projData;
+ D_subprojData.ptr = (char*)D_projData.ptr + iAngle * D_projData.pitch;
+
+ ret = ConeFP_Array_internal(D_subprojData,
+ dims, iEndAngle - iAngle, angles + iAngle,
+ fOutputScale);
+ if (!ret)
+ break;
+ }
cudaFreeArray(cuArray);