#include <stdio.h>

#include "saft.hpp"

#define DebugSosVoxelX 18
#define DebugSosVoxelY 18
#define DebugSosVoxelZ 3
// Surfaces fuer Emitter - SosPathsTables
extern surface<void, cudaSurfaceType3D> outSurfRefTableVoxelToEmPathSosBoth;

// Surfaces fuer Emitter - SosPathsTables
// surface <void, cudaSurfaceType3D>
// outSurfRefTableVoxelToReceiverPathSosSumTest;
extern surface<void, cudaSurfaceType3D> outSurfRefTableVoxelToRecPathSosBoth0;
extern surface<void, cudaSurfaceType3D> outSurfRefTableVoxelToRecPathSosBoth1;
extern surface<void, cudaSurfaceType3D> outSurfRefTableVoxelToRecPathSosBoth2;

// Surfaces fuer AscanIndex
extern surface<void, cudaSurfaceType3D> outSurfRefAscanIndexFloat0;
extern surface<void, cudaSurfaceType3D> outSurfRefAscanIndexFloat1;
extern surface<void, cudaSurfaceType3D> outSurfRefAscanIndexFloat2;
extern surface<void, cudaSurfaceType3D> outSurfRefAscanIndexFloat3;

// Texturmemory fuer Emitter - SosPathsTables

extern texture<float4, cudaTextureType3D, cudaReadModeElementType> texTableVoxelToEmitterPathSosBoth_preprocess;
// Texturmemory fuer Receiver - SosPathsTables

extern texture<float4, cudaTextureType3D, cudaReadModeElementType> texTableVoxelToReceiverPathSosBoth0_preprocess;
extern texture<float4, cudaTextureType3D, cudaReadModeElementType> texTableVoxelToReceiverPathSosBoth1_preprocess;
extern texture<float4, cudaTextureType3D, cudaReadModeElementType> texTableVoxelToReceiverPathSosBoth2_preprocess;

__global__ void precalculateAverageSpeedOfSoundKernel(cudaArray *deviceSosAttFieldCuArray,
                                                      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,
                                                      float debugModeParameter);

__global__ void precalculateAscanIndex_usePathsKernel(int ascanIndexBatchOffset, 
                                                      int aScanWindowSize,
                                                      cudaArray *deviceSosAttFieldCuArray, 
                                                      int currentSpeedOfSoundZLayer, 
                                                      int maxFeasibleSosZLayerCount, 
                                                      int maxSoSReceiverArrayForTexture,
                                                      unsigned short const *deviceEmitterIndex_block,
                                                      unsigned short const *deviceReceiverIndex_block,
                                                      int TableAscanIndexAllocationCount, 
                                                      int maxAscanIndexArraysInTexture, 
                                                      cudaArray **deviceTextureAscanIndexFloatCuArray, int3 SOSGrid_XYZ, float3 sosOffset, float3 regionOfInterestOffset, float IMAGE_RESOLUTION,
                                                      float SOS_RESOLUTION, float debugMode, float debugModeParameter, int *deviceSAFT_VARIANT);

__global__ void precalculateAscanIndex_usePathsKernel_SOS(int ascanIndexBatchOffset,
                                                          int aScanWindowSize, 
                                                          cudaArray *deviceSosAttFieldCuArray,
                                                          int currentSpeedOfSoundZLayer, 
                                                          int maxFeasibleSosZLayerCount,
                                                          int maxSoSReceiverArrayForTexture,
                                                          unsigned short const *deviceEmitterIndex_block,
                                                          unsigned short const *deviceReceiverIndex_block, 
                                                          int TableAscanIndexAllocationCount,
                                                          int maxAscanIndexArraysInTexture, 
                                                          cudaArray **deviceTextureAscanIndexFloatCuArray, int3 SOSGrid_XYZ, float3 sosOffset, float3 regionOfInterestOffset, float IMAGE_RESOLUTION,
                                                          float SOS_RESOLUTION, float debugMode, float debugModeParameter, int *deviceSAFT_VARIANT);

__global__ void precalculateAscanIndex_usePathsKernel_SOS_ATT(int ascanIndexBatchOffset,
                                                              int aScanWindowSize,
                                                              cudaArray *deviceSosAttFieldCuArray,
                                                              int currentSpeedOfSoundZLayer,
                                                              int maxFeasibleSosZLayerCount, 
                                                              int maxSoSReceiverArrayForTexture,
                                                              unsigned short const *deviceEmitterIndex_block, 
                                                              unsigned short const *deviceReceiverIndex_block,
                                                              int TableAscanIndexAllocationCount,
                                                              int maxAscanIndexArraysInTexture,
                                                              cudaArray **deviceTextureAscanIndexFloatCuArray,

                                                              int3 SOSGrid_XYZ, float3 sosOffset, float3 regionOfInterestOffset, float IMAGE_RESOLUTION, float SOS_RESOLUTION, float debugMode,
                                                              float debugModeParameter, int *deviceSAFT_VARIANT);

__global__ void fillCuArrayKernel(
    float useValue,
    cudaArray **deviceTextureAscanIndexFloatCuArray,  
    int maxAscanIndexArraysInTexture, int TableAscanIndexAllocationCount,
    int maxFeasibleSosZLayerCount,  
    bool ATTMode_3DVolume,
    float debugMode, float debugModeParameter);