142 lines
4.0 KiB
C++
142 lines
4.0 KiB
C++
#ifndef vtk_m_worklet_zfp_encode1_h
|
|
#define vtk_m_worklet_zfp_encode1_h
|
|
|
|
#include <vtkm/Types.h>
|
|
#include <vtkm/internal/ExportMacros.h>
|
|
|
|
#include <vtkm/worklet/WorkletMapField.h>
|
|
#include <vtkm/worklet/zfp/ZFPBlockWriter.h>
|
|
#include <vtkm/worklet/zfp/ZFPEncode.h>
|
|
#include <vtkm/worklet/zfp/ZFPFunctions.h>
|
|
#include <vtkm/worklet/zfp/ZFPStructs.h>
|
|
#include <vtkm/worklet/zfp/ZFPTypeInfo.h>
|
|
|
|
|
|
|
|
namespace vtkm
|
|
{
|
|
namespace worklet
|
|
{
|
|
namespace zfp
|
|
{
|
|
|
|
template <typename Scalar, typename PortalType>
|
|
VTKM_EXEC inline void GatherPartial1(Scalar* q,
|
|
const PortalType& scalars,
|
|
vtkm::Id offset,
|
|
int nx,
|
|
int sx)
|
|
{
|
|
vtkm::Id x;
|
|
for (x = 0; x < nx; x++, offset += sx)
|
|
q[x] = scalars.Get(offset);
|
|
PadBlock(q, nx, 1);
|
|
}
|
|
|
|
template <typename Scalar, typename PortalType>
|
|
VTKM_EXEC inline void Gather1(Scalar* fblock, const PortalType& scalars, vtkm::Id offset, int sx)
|
|
{
|
|
vtkm::Id counter = 0;
|
|
|
|
for (vtkm::Id x = 0; x < 4; x++, offset += sx)
|
|
{
|
|
fblock[counter] = scalars.Get(offset);
|
|
counter++;
|
|
}
|
|
}
|
|
|
|
struct Encode1 : public vtkm::worklet::WorkletMapField
|
|
{
|
|
protected:
|
|
vtkm::Id Dims; // field dims
|
|
vtkm::Id PaddedDims; // dims padded to a multiple of zfp block size
|
|
vtkm::Id ZFPDims; // zfp block dims
|
|
vtkm::UInt32 MaxBits; // bits per zfp block
|
|
|
|
public:
|
|
Encode1(const vtkm::Id dims, const vtkm::Id paddedDims, const vtkm::UInt32 maxbits)
|
|
: Dims(dims)
|
|
, PaddedDims(paddedDims)
|
|
, MaxBits(maxbits)
|
|
{
|
|
ZFPDims = PaddedDims / 4;
|
|
}
|
|
using ControlSignature = void(FieldIn<>, WholeArrayIn<>, AtomicArrayInOut<> bitstream);
|
|
using ExecutionSignature = void(_1, _2, _3);
|
|
|
|
template <class InputScalarPortal, typename BitstreamPortal>
|
|
VTKM_EXEC void operator()(const vtkm::Id blockIdx,
|
|
const InputScalarPortal& scalars,
|
|
BitstreamPortal& stream) const
|
|
{
|
|
using Scalar = typename InputScalarPortal::ValueType;
|
|
|
|
// typedef unsigned long long int ull;
|
|
// typedef long long int ll;
|
|
// const ull blockId = blockIdx.x +
|
|
// blockIdx.y * gridDim.x +
|
|
// gridDim.x * gridDim.y * blockIdx.z;
|
|
|
|
// // each thread gets a block so the block index is
|
|
// // the global thread index
|
|
// const uint block_idx = blockId * blockDim.x + threadIdx.x;
|
|
|
|
// if(block_idx >= tot_blocks)
|
|
// {
|
|
// // we can't launch the exact number of blocks
|
|
// // so just exit if this isn't real
|
|
// return;
|
|
// }
|
|
|
|
// uint2 block_dims;
|
|
// block_dims.x = padded_dims.x >> 2;
|
|
// block_dims.y = padded_dims.y >> 2;
|
|
|
|
// // logical pos in 3d array
|
|
// uint2 block;
|
|
// block.x = (block_idx % block_dims.x) * 4;
|
|
// block.y = ((block_idx/ block_dims.x) % block_dims.y) * 4;
|
|
// const ll offset = (ll)block.x * stride.x + (ll)block.y * stride.y;
|
|
|
|
vtkm::Id zfpBlock;
|
|
zfpBlock = blockIdx % ZFPDims;
|
|
vtkm::Id logicalStart = zfpBlock * vtkm::Id(4);
|
|
|
|
constexpr vtkm::Int32 BlockSize = 4;
|
|
Scalar fblock[BlockSize];
|
|
|
|
// bool partial = false;
|
|
// if(block.x + 4 > dims.x) partial = true;
|
|
// if(block.y + 4 > dims.y) partial = true;
|
|
|
|
bool partial = false;
|
|
if (logicalStart + 4 > Dims)
|
|
partial = true;
|
|
|
|
if (partial)
|
|
{
|
|
const vtkm::Int32 nx =
|
|
logicalStart + 4 > Dims ? vtkm::Int32(Dims - logicalStart) : vtkm::Int32(4);
|
|
GatherPartial1(fblock, scalars, logicalStart, nx, 1);
|
|
}
|
|
else
|
|
{
|
|
Gather1(fblock, scalars, logicalStart, 1);
|
|
}
|
|
|
|
for (int i = 0; i < 16; ++i)
|
|
{
|
|
std::cout << " " << fblock[i];
|
|
}
|
|
std::cout << "\n";
|
|
|
|
//zfp_encode_block<Scalar, ZFP_2D_BLOCK_SIZE>(fblock, maxbits, block_idx, stream);
|
|
zfp::ZFPBlockEncoder<BlockSize, Scalar, BitstreamPortal> encoder;
|
|
encoder.encode(fblock, MaxBits, vtkm::UInt32(blockIdx), stream);
|
|
}
|
|
};
|
|
}
|
|
}
|
|
}
|
|
#endif
|