diff options
| author | Willem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl> | 2021-12-02 11:52:10 +0100 |
|---|---|---|
| committer | Willem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl> | 2021-12-06 11:01:36 +0100 |
| commit | 33f52988134d11096b69352671c54045a30a82d4 (patch) | |
| tree | 92d810de0471deb5f78ac7fa87213c9f33caafb2 /cuda/3d | |
| parent | 86615d4161b050fbf3335e30ae85801aa1cefe92 (diff) | |
| download | astra-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/3d')
| -rw-r--r-- | cuda/3d/cone_fp.cu | 51 | ||||
| -rw-r--r-- | cuda/3d/par3d_fp.cu | 49 |
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 |