diff options
Diffstat (limited to 'cuda/3d/cone_fp.cu')
| -rw-r--r-- | cuda/3d/cone_fp.cu | 51 |
1 files changed, 30 insertions, 21 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 |