summaryrefslogtreecommitdiffstats
path: root/cuda
diff options
context:
space:
mode:
authorWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2021-12-02 11:45:41 +0100
committerWillem Jan Palenstijn <Willem.Jan.Palenstijn@cwi.nl>2021-12-02 11:52:38 +0100
commit0070b8bafe1365b97ed922bd0ede84de7edd7dc7 (patch)
tree353068bf6e2903805d9c911359f0ed7698c7cb95 /cuda
parentc686f9ac537f7ebf66ea7d1fff7df19498c09525 (diff)
downloadastra-0070b8bafe1365b97ed922bd0ede84de7edd7dc7.tar.gz
astra-0070b8bafe1365b97ed922bd0ede84de7edd7dc7.tar.bz2
astra-0070b8bafe1365b97ed922bd0ede84de7edd7dc7.tar.xz
astra-0070b8bafe1365b97ed922bd0ede84de7edd7dc7.zip
Fix potential invalid memory access in par3d_fp with many angles
Diffstat (limited to 'cuda')
-rw-r--r--cuda/3d/par3d_fp.cu4
1 files changed, 2 insertions, 2 deletions
diff --git a/cuda/3d/par3d_fp.cu b/cuda/3d/par3d_fp.cu
index b2178ec..9475897 100644
--- a/cuda/3d/par3d_fp.cu
+++ b/cuda/3d/par3d_fp.cu
@@ -401,7 +401,7 @@ bool Par3DFP_Array_internal(cudaPitchedPtr D_projData,
const SProjectorParams3D& params)
{
// 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 < angleCount; ++i) tmp[i] = angles[i].f##name ; cudaMemcpyToSymbol(gC_##name, tmp, angleCount*sizeof(float), 0, cudaMemcpyHostToDevice); } while (0)
@@ -468,7 +468,7 @@ bool Par3DFP_Array_internal(cudaPitchedPtr D_projData,
for (unsigned int a = 0; a <= angleCount; ++a) {
int dir = -1;
- if (a != dims.iProjAngles) {
+ if (a != angleCount) {
float dX = fabsf(angles[a].fRayX);
float dY = fabsf(angles[a].fRayY);
float dZ = fabsf(angles[a].fRayZ);