diff --git a/SAFT_TOFI/src/kernel/precalculateSpeedOfSoundKernel.cu b/SAFT_TOFI/src/kernel/precalculateSpeedOfSoundKernel.cu index 57703bb..496d91d 100644 --- a/SAFT_TOFI/src/kernel/precalculateSpeedOfSoundKernel.cu +++ b/SAFT_TOFI/src/kernel/precalculateSpeedOfSoundKernel.cu @@ -45,21 +45,15 @@ texture texTableVoxelToRecei texture texTableVoxelToReceiverPathSosBoth1_preprocess; texture texTableVoxelToReceiverPathSosBoth2_preprocess; -__global__ void precalculateAverageSpeedOfSoundKernel(cudaArray *deviceSosAttFieldCuArray, ///< CuArray fuer SosAttFieldTextur - int firstZLayer, ///< First z-layer in the speed of sound grid the - ///< pre-calculation is performed for. - int sosZLayerCount, ///< Number of z-layers in the speed of sound grid the - ///< pre-calculation is performed for. - int geometry, ///< emitters=0 or receivers=1. - int geometryElementCount, ///< Number of elements in the geometry array. - int maxSoSReceiverArrayForTexture, ///< max amount of elements in the - ///< receiver CUDA array. - float *deviceVoxelCountOutputFloat, ///< fuer Count im Floatformat gedacht - ///< fuer Texturmemory. - float *speedOfSoundSumOutput, ///< fuer SoS im Floatformat gedacht fuer - ///< Texturmemory. - int3 SOSGrid_XYZ, float3 sosOffset, float3 regionOfInterestOffset, float IMAGE_RESOLUTION, float SOS_RESOLUTION, float debugMode, - float debugModeParameter) +__global__ void precalculateAverageSpeedOfSoundKernel(int firstZLayer, ///< First z-layer in the speed of sound grid the + ///< pre-calculation is performed for. + int sosZLayerCount, ///< Number of z-layers in the speed of sound grid the + ///< pre-calculation is performed for. + int geometry, ///< emitters=0 or receivers=1. + int geometryElementCount, ///< Number of elements in the geometry array. + int maxSoSReceiverArrayForTexture, ///< max amount of elements in the + ///< receiver CUDA array. + float3 sosOffset, float3 regionOfInterestOffset, float IMAGE_RESOLUTION, float SOS_RESOLUTION, float debugMode, float debugModeParameter) { dim3 SosVoxel(threadIdx.x, // SoS-Voxel X ? Threads fangen bei 0 an blockIdx.x, // SoS-Voxel Y @@ -123,12 +117,10 @@ __global__ void precalculateAverageSpeedOfSoundKernel(cudaArray *deviceSosAttFie // Bestimmen der SoS-Koordinaten fuer die Sender/Empfuenger-Koordinaten determineSpeedOfSoundFieldVoxelFloat(currentGeometry, SosGeometryVoxelFloat, sosOffset, - SOS_RESOLUTION_FACTOR); // SoSVoxel von E/R bestimmen // - // currentGeometry --> SosGeometryVoxel Float - performRayTracedSpeedAdditionTexture(voxelCount, averageSpeed, totalAttenuation, SosGeometryVoxelFloat, SosVoxel, deviceSosAttFieldCuArray, SOSGrid_XYZ, sosOffset, SOS_RESOLUTION, - IMAGE_RESOLUTION, regionOfInterestOffset, - geometry); // SosGeometryVoxelFloat im Floatformat, SoSVoxel als - // Integer + SOS_RESOLUTION_FACTOR); // SoSVoxel von E/R bestimmen // + // currentGeometry --> SosGeometryVoxel Float + performRayTracedSpeedAdditionTexture(voxelCount, averageSpeed, totalAttenuation, SosGeometryVoxelFloat, SosVoxel, sosOffset); // SosGeometryVoxelFloat im Floatformat, + // SoSVoxel als Integer if (geometry == 0) // Emitter { diff --git a/SAFT_TOFI/src/kernel/precalculateSpeedOfSoundKernel.cuh b/SAFT_TOFI/src/kernel/precalculateSpeedOfSoundKernel.cuh index 773ea0c..279c45d 100644 --- a/SAFT_TOFI/src/kernel/precalculateSpeedOfSoundKernel.cuh +++ b/SAFT_TOFI/src/kernel/precalculateSpeedOfSoundKernel.cuh @@ -30,15 +30,12 @@ extern texture texTableVoxel extern texture texTableVoxelToReceiverPathSosBoth1_preprocess; extern texture texTableVoxelToReceiverPathSosBoth2_preprocess; -__global__ void precalculateAverageSpeedOfSoundKernel(cudaArray *deviceSosAttFieldCuArray, - int firstZLayer, +__global__ void precalculateAverageSpeedOfSoundKernel(int firstZLayer, int sosZLayerCount, int geometry, int geometryElementCount, int maxSoSReceiverArrayForTexture, - float *deviceVoxelCountOutputFloat, - float *speedOfSoundSumOutput, - int3 SOSGrid_XYZ, float3 sosOffset, float3 regionOfInterestOffset, float IMAGE_RESOLUTION, float SOS_RESOLUTION, float debugMode, + float3 sosOffset, float3 regionOfInterestOffset, float IMAGE_RESOLUTION, float SOS_RESOLUTION, float debugMode, float debugModeParameter); __global__ void precalculateAscanIndex_usePathsKernel(int ascanIndexBatchOffset, diff --git a/SAFT_TOFI/src/kernel/rayTracing.cuh b/SAFT_TOFI/src/kernel/rayTracing.cuh index 2330239..bf57c68 100644 --- a/SAFT_TOFI/src/kernel/rayTracing.cuh +++ b/SAFT_TOFI/src/kernel/rayTracing.cuh @@ -45,9 +45,7 @@ __device__ __forceinline__ void processRayTracedVoxelTextureSosAtt( float const *currentVoxelFloat, ///< SOS-coordinates of current Voxel for Bresenham in float[3]. float &voxelCount, ///< This argument is written to. Number of voxels in the path so far. // TODO: not necessary here to sum up because it is already known from Bresenham length float &totalSpeed, ///< This argument is written to. Sum Up total SOS - float &totalAttenuation, ///< This argument is written to. Sum Up total Attenuation - cudaArray *deviceSosAttFieldCuArray, ///< Pointer to cudaArray for SoSATTFieldData // TODO: not necessary if access is used with texture memory - int3 const &SOSGrid_XYZ ///< Size of SOS-Grid in XYZ // TODO: only necessary for boundary-check. + float &totalAttenuation ///< This argument is written to. Sum Up total Attenuation ) { @@ -59,19 +57,11 @@ __device__ __forceinline__ void processRayTracedVoxelTextureSosAtt( __device__ __forceinline__ void performRayTracedSpeedAdditionTexture(float &voxelCount, ///< This argument is written to. Number of voxels within the path traced. - float &totalSpeed, ///< This argument is written to. Sum of the speed of sound samples in the path traced. float &totalAttenuation, - float3 const &point1f, ///< Vector array describing the Voxelcoordinates of emitters or receivers. dim3 const &point2, ///< Vector array describing the Voxelcoordinates of Voxels. - - cudaArray *deviceSpeedOfSoundFieldCuArray, ///< CuArray fuer SOSFieldTextur - - int3 const &SOSGrid_XYZ, ///< Size of SOS-Grid in XYZ - float3 &sosOffset, float &SOS_RESOLUTION, float &IMAGE_RESOLUTION, float3 regionOfInterestOffset, - int geometry ///< emitters=0 or receivers=1. -) + float3 &sosOffset) { voxelCount = 0.0f; @@ -156,7 +146,7 @@ __device__ __forceinline__ void performRayTracedSpeedAdditionTexture(float &voxe int j = 0; for (j = fastDirectionSteps; j > 0; j--) //(Alle Punkte innerhalb der Schleife berechnen) { - processRayTracedVoxelTextureSosAtt(pathPoint, voxelCount, totalSpeed, totalAttenuation, deviceSpeedOfSoundFieldCuArray, SOSGrid_XYZ); + processRayTracedVoxelTextureSosAtt(pathPoint, voxelCount, totalSpeed, totalAttenuation); pathPoint[greatestDistanceDim] = pathPoint[greatestDistanceDim] + m_XYZ[greatestDistanceDim]; pathPoint[slowDim1] = pathPoint[slowDim1] + m_XYZ[slowDim1]; pathPoint[slowDim2] = pathPoint[slowDim2] + m_XYZ[slowDim2]; diff --git a/SAFT_TOFI/src/kernel/saftPrivate.cu b/SAFT_TOFI/src/kernel/saftPrivate.cu index 1b67a8a..6fa33c5 100644 --- a/SAFT_TOFI/src/kernel/saftPrivate.cu +++ b/SAFT_TOFI/src/kernel/saftPrivate.cu @@ -228,8 +228,7 @@ void SAFTHandler::precalculateAscanIndex_usePaths(int ascanIndex_i, int aScanWin CUDA_CHECK(cudaUnbindTexture(&texTableVoxelToReceiverPathSosBoth2_preprocess)); } -void SAFTHandler::precalculateAverageSpeedOfSound(int firstZLayer, int sosZLayerCount, int deviceListGeometry, int geometryElementCount, float *deviceVoxelCountOutputFloat, - float *deviceSpeedOfSoundSumOutput) +void SAFTHandler::precalculateAverageSpeedOfSound(int firstZLayer, int sosZLayerCount, int deviceListGeometry, int geometryElementCount) { dim3 threadsPerBlock(SOSGrid_XYZ.x, 1, 1); // max. 512 oder 1024 Threads werden vorgegeben und @@ -284,14 +283,9 @@ void SAFTHandler::precalculateAverageSpeedOfSound(int firstZLayer, int sosZLayer } } - precalculateAverageSpeedOfSoundKernel<<>>(deviceSosAttFieldCuArray, firstZLayer, sosZLayerCount, deviceListGeometry, geometryElementCount, - maxSoSReceiverArrayForTexture, // maximale Anzahl an Receivern in einem - // CUDA Array - - // deviceVoxelCountOutput, - deviceVoxelCountOutputFloat, deviceSpeedOfSoundSumOutput, - // regionOfInterestOffset, - SOSGrid_XYZ, sosOffset, regionOfInterestOffset, IMAGE_RESOLUTION, SOS_RESOLUTION, debugMode, debugModeParameter); + precalculateAverageSpeedOfSoundKernel<<>>(firstZLayer, sosZLayerCount, deviceListGeometry, geometryElementCount, + maxSoSReceiverArrayForTexture, + sosOffset, regionOfInterestOffset, IMAGE_RESOLUTION, SOS_RESOLUTION, debugMode, debugModeParameter); CUDA_CHECK(cudaGetLastError()); CUDA_CHECK(cudaUnbindTexture(&texRefSosAttField)); diff --git a/SAFT_TOFI/src/processAScans.cpp b/SAFT_TOFI/src/processAScans.cpp index f3abe10..741b549 100644 --- a/SAFT_TOFI/src/processAScans.cpp +++ b/SAFT_TOFI/src/processAScans.cpp @@ -130,12 +130,11 @@ void SAFTHandler::performCoreReconstruction() // SoS-Pfade für Z-Layer vorberechnen ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ // Emitter ----------------------------------------------------------------------- Emitter - precalculateAverageSpeedOfSound(currentSpeedOfSoundZLayer, maxFeasibleSosZLayerCount, 0, emitter_list_Size, deviceTableVoxelToEmitterPathCountFloat, deviceTableVoxelToEmitterPathSosSum); + precalculateAverageSpeedOfSound(currentSpeedOfSoundZLayer, maxFeasibleSosZLayerCount, 0, emitter_list_Size); // Receiver ----------------------------------------------------------------------- Receiver - precalculateAverageSpeedOfSound(currentSpeedOfSoundZLayer, maxFeasibleSosZLayerCount, 1, receiver_list_Size, deviceTableVoxelToReceiverPathCountFloat, - deviceTableVoxelToReceiverPathSosSum); + precalculateAverageSpeedOfSound(currentSpeedOfSoundZLayer, maxFeasibleSosZLayerCount, 1, receiver_list_Size); } // Go overall A-Scanblocks ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ // PrecalAndPerformSAFTAllAscans - Berechne Performanz fuer Durchlauf von allen Ascans ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ diff --git a/SAFT_TOFI/src/saft.hpp b/SAFT_TOFI/src/saft.hpp index a9f361b..8b620b6 100644 --- a/SAFT_TOFI/src/saft.hpp +++ b/SAFT_TOFI/src/saft.hpp @@ -364,7 +364,7 @@ class SAFTHandler // Pre-calculation kernels //------------------------------------------------------------------------ - void precalculateAverageSpeedOfSound(int firstZLayer, int sosZLayerCount, int deviceGeometry, int geometryElementCount, float *deviceVoxelCountOutputFloat, float *deviceSpeedOfSoundSumOutput); + void precalculateAverageSpeedOfSound(int firstZLayer, int sosZLayerCount, int deviceGeometry, int geometryElementCount); void precalculateAscanIndex(int currentSpeedOfSoundZLayer, ///< First z-layer in the speed of sound grid the pre-calculation is performed for. int maxFeasibleSosZLayerCount ///< Number of z-layers in the speed of sound grid the pre-calculation is performed for.