samplingTest.cpp

// Copyright (c) 2012-2017 VideoStitch SAS
// Copyright (c) 2018 stitchEm

#include "common/testing.hpp"
#include "common/util.hpp"
#include "common/shiftBuffers.hpp"

#include "libvideostitch/logging.hpp"
#include "libvideostitch/profile.hpp"
#include "libvideostitch/gpu_device.hpp"
#include <gpu/stream.hpp>
#include <gpu/image/sampling.hpp>

#include <stdint.h>

namespace VideoStitch {
namespace Testing {
void testSubsample22(const unsigned char* in, const unsigned char* expOut, unsigned inWidth, unsigned inHeight) {
  DeviceBuffer<unsigned char> devInBuffer(inWidth, inHeight);
  devInBuffer.fillData(in);
  const unsigned outWidth = (inWidth + 1) / 2;
  const unsigned outHeight = (inHeight + 1) / 2;
  DeviceBuffer<unsigned char> devOutBuffer(outWidth, outHeight);
  Image::subsample22(devOutBuffer.gpuBuf(), devInBuffer.gpuBufConst(), inWidth, inHeight, GPU::Stream::getDefault());
  GPU::Stream::getDefault().synchronize();
  std::vector<unsigned char> actualOut;
  devOutBuffer.readback(actualOut);

  for (unsigned y = 0; y < outHeight; ++y) {
    for (unsigned x = 0; x < outWidth; ++x) {
      ENSURE_EQ(expOut[y * outWidth + x], actualOut[y * outWidth + x]);
    }
  }
}
/*
void testSubsample22Mask(const uint32_t *in, const uint32_t* inMask, const uint32_t *expOut, const uint32_t* expOutMask,
unsigned inWidth, unsigned inHeight) { DeviceBuffer<uint32_t> devInBuffer(inWidth, inHeight); devInBuffer.fillData(in);

  DeviceBuffer<uint32_t> devInMaskBuffer(inWidth, inHeight);
  devInMaskBuffer.fillData(inMask);

  const unsigned outWidth = (inWidth + 1) / 2;
  const unsigned outHeight = (inHeight + 1) / 2;
  DeviceBuffer<uint32_t> devOutBuffer(outWidth, outHeight);
  DeviceBuffer<uint32_t> devOutMaskBuffer(outWidth, outHeight);

  Image::subsample22Mask<uint32_t>(devOutBuffer.gpuBuf(), devOutMaskBuffer.gpuBuf(), devInBuffer.gpuBufConst(),
devInMaskBuffer.gpuBufConst(), inWidth, inHeight, 16, GPU::Stream::getDefault());
  GPU::Stream::getDefault().synchronize();
  std::vector<uint32_t> actualOut;
  devOutBuffer.readback(actualOut);
  std::vector<uint32_t> actualOutMask;
  devOutMaskBuffer.readback(actualOutMask);

  for (unsigned y = 0; y < outHeight; ++y) {
    for (unsigned x = 0; x < outWidth; ++x) {
      ENSURE_EQ(expOut[y * outWidth + x], actualOut[y * outWidth + x]);
      ENSURE_EQ(expOutMask[y * outWidth + x], actualOutMask[y * outWidth + x]);
    }
  }
}

void testSubsample22Nearest(const uint32_t* in, const uint32_t* expOut, unsigned inWidth, unsigned inHeight) {
  DeviceBuffer<uint32_t> devInBuffer(inWidth, inHeight);
  devInBuffer.fillData(in);
  const unsigned outWidth = (inWidth + 1) / 2;
  const unsigned outHeight = (inHeight + 1) / 2;
  DeviceBuffer<uint32_t> devOutBuffer(outWidth, outHeight);
  Image::subsample22Nearest(devOutBuffer.gpuBuf(), devInBuffer.gpuBufConst(), inWidth, inHeight, 16,
GPU::Stream::getDefault()); GPU::Stream::getDefault().synchronize(); std::vector<uint32_t> actualOut;
  devOutBuffer.readback(actualOut);

  for (unsigned y = 0; y < outHeight; ++y) {
    for (unsigned x = 0; x < outWidth; ++x) {
      ENSURE_EQ(expOut[y * outWidth + x], actualOut[y * outWidth + x]);
    }
  }
}

void testSubsampleMask22(const unsigned char* in, const unsigned char* expOut, unsigned inWidth, unsigned inHeight) {
  DeviceBuffer<unsigned char> devInBuffer(inWidth, inHeight);
  devInBuffer.fillData(in);
  const unsigned outWidth = (inWidth + 1) / 2;
  const unsigned outHeight = (inHeight + 1) / 2;
  DeviceBuffer<unsigned char> devOutBuffer(outWidth, outHeight);
  Image::subsampleMask22(devOutBuffer.gpuBuf(), devInBuffer.gpuBuf(), inWidth, inHeight, 16, GPU::Stream::getDefault());
  GPU::Stream::getDefault().synchronize();
  std::vector<unsigned char> actualOut;
  devOutBuffer.readback(actualOut);

  for (unsigned y = 0; y < outHeight; ++y) {
    for (unsigned x = 0; x < outWidth; ++x) {
      ENSURE_EQ((int)expOut[y * outWidth + x], (int)actualOut[y * outWidth + x]);
    }
  }
}
*/
void testSubsample22RGBA(const uint32_t* in, const uint32_t* expOut, unsigned inWidth, unsigned inHeight) {
  DeviceBuffer<uint32_t> devInBuffer(inWidth, inHeight);
  devInBuffer.fillData(in);
  const unsigned outWidth = (inWidth + 1) / 2;
  const unsigned outHeight = (inHeight + 1) / 2;
  DeviceBuffer<uint32_t> devOutBuffer(outWidth, outHeight);
  Image::subsample22RGBA(devOutBuffer.gpuBuf(), devInBuffer.gpuBuf(), inWidth, inHeight, GPU::Stream::getDefault());
  GPU::Stream::getDefault().synchronize();
  std::vector<uint32_t> actualOut;
  devOutBuffer.readback(actualOut);

  ENSURE_RGBA8888_ARRAY_EQ(expOut, actualOut.data(), outWidth, outHeight);
}

void testUpsample22(const uint32_t* in, const uint32_t* expOut, unsigned outWidth, unsigned outHeight) {
  const unsigned inWidth = (outWidth + 1) / 2;
  const unsigned inHeight = (outHeight + 1) / 2;
  DeviceBuffer<uint32_t> devInBuffer(inWidth, inHeight);
  devInBuffer.fillData(in);
  DeviceBuffer<uint32_t> devOutBuffer(outWidth, outHeight);
  Image::upsample22(devOutBuffer.gpuBuf(), devInBuffer.gpuBufConst(), outWidth, outHeight, false,
                    GPU::Stream::getDefault());
  GPU::Stream::getDefault().synchronize();
  std::vector<uint32_t> actualOut;
  devOutBuffer.readback(actualOut);

  for (unsigned y = 0; y < outHeight; ++y) {
    for (unsigned x = 0; x < outWidth; ++x) {
      ENSURE_EQ(expOut[y * outWidth + x], actualOut[y * outWidth + x]);
    }
  }
}

void testUpsample22RGBA210(const uint32_t* inRGBA, const uint32_t* expOutRGBA, unsigned outWidth, unsigned outHeight) {
  const unsigned inWidth = (outWidth + 1) / 2;
  const unsigned inHeight = (outHeight + 1) / 2;
  DeviceBuffer<uint32_t> devInBuffer(inWidth, inHeight);
  devInBuffer.fillData(inRGBA);
  DeviceBuffer<uint32_t> devOutBuffer(outWidth, outHeight);
  Image::upsample22RGBA210(devOutBuffer.gpuBuf(), devInBuffer.gpuBuf(), outWidth, outHeight, false,
                           GPU::Stream::getDefault());
  GPU::Stream::getDefault().synchronize();
  std::vector<uint32_t> actualOut;
  devOutBuffer.readback(actualOut);

  ENSURE_RGBA210_ARRAY_EQ(expOutRGBA, actualOut.data(), outWidth, outHeight);
}

/**
 * Upsampling then subsampling is idempotent.
 * upsampling the result gives back the same result.
 */
void testRoundTripsample(const unsigned char* in, unsigned upWidth, unsigned upHeight, unsigned blockSize) {
  // Convert buffer to RGBA
  std::vector<uint32_t> upRGB;
  for (unsigned i = 0; i < upWidth * upHeight; ++i) {
    upRGB.push_back(Image::RGBA::pack(in[3 * i + 0], in[3 * i + 1], in[3 * i + 2], 0xff));
  }

  unsigned downWidth = (upWidth + 1) / 2;
  unsigned downHeight = (upHeight + 1) / 2;

  DeviceBuffer<uint32_t> devUpBuffer(upWidth, upHeight);
  DeviceBuffer<uint32_t> devUpBuffer2(upWidth, upHeight);
  devUpBuffer.fill(upRGB);
  DeviceBuffer<uint32_t> devDownBuffer(downWidth, downHeight);
  DeviceBuffer<uint32_t> devDownBuffer2(downWidth, downHeight);
  Image::subsample22RGBA(devDownBuffer.gpuBuf(), devUpBuffer.gpuBufConst(), upWidth, upHeight,
                         GPU::Stream::getDefault());
  GPU::Stream::getDefault().synchronize();

  // Round trips
  Image::upsample22RGBA(devUpBuffer.gpuBuf(), devDownBuffer.gpuBuf(), upWidth, upHeight, false,
                        GPU::Stream::getDefault());
  GPU::Stream::getDefault().synchronize();
  Image::subsample22RGBA(devDownBuffer2.gpuBuf(), devUpBuffer.gpuBufConst(), upWidth, upHeight,
                         GPU::Stream::getDefault());
  GPU::Stream::getDefault().synchronize();
  Image::upsample22RGBA(devUpBuffer2.gpuBuf(), devDownBuffer2.gpuBuf(), upWidth, upHeight, false,
                        GPU::Stream::getDefault());
  GPU::Stream::getDefault().synchronize();

  std::vector<uint32_t> hostUpBuffer;
  devUpBuffer.readback(hostUpBuffer);
  std::vector<uint32_t> hostUpBuffer2;
  devUpBuffer2.readback(hostUpBuffer2);
  std::vector<uint32_t> hostDownBuffer;
  devDownBuffer.readback(hostDownBuffer);
  std::vector<uint32_t> hostDownBuffer2;
  devDownBuffer2.readback(hostDownBuffer2);

  for (unsigned y = 0; y < downHeight; ++y) {
    for (unsigned x = 0; x < downWidth; ++x) {
      ENSURE_EQ(hostDownBuffer[y * downWidth + x], hostDownBuffer2[y * downWidth + x]);
    }
  }

  for (unsigned y = 0; y < upHeight; ++y) {
    for (unsigned x = 0; x < upWidth; ++x) {
      ENSURE_EQ(hostUpBuffer[y * upWidth + x], hostUpBuffer2[y * upWidth + x]);
    }
  }
}

/**
 * Make sure that wrapping upsampling works:
 * shifting then upsampling should be the same as subsampling then shifting.
 */
void testWrappingUpsample22RGBA210(unsigned upWidth, unsigned upHeight, unsigned blockSize, unsigned shift) {
  unsigned downWidth = (upWidth + 1) / 2;
  unsigned downHeight = (upHeight + 1) / 2;
  // Convert buffer to RGBA
  std::vector<uint32_t> downRGBA;
  int v = 7;
  for (unsigned i = 0; i < downWidth * downHeight; ++i) {
    int32_t r = v;
    v = (v * 13) % 255;
    int32_t g = v;
    v = (v * 13) % 255;
    int32_t b = v;
    v = (v * 13) % 255;
    downRGBA.push_back(Image::RGB210::pack(r, g, b, 0xff));
  }

  DeviceBuffer<uint32_t> devDownBuffer(downWidth, downHeight);
  devDownBuffer.fill(downRGBA);
  DeviceBuffer<uint32_t> devUpBuffer(upWidth, upHeight);

  // upsample unshifted
  Image::upsample22RGBA210(devUpBuffer.gpuBuf(), devDownBuffer.gpuBuf(), upWidth, upHeight, true,
                           GPU::Stream::getDefault());
  cudaStreamSynchronize(0);

  std::vector<uint32_t> hostUpBuffer;
  devUpBuffer.readback(hostUpBuffer);

  // shift the input
  shiftDevLeft(devDownBuffer.gpuBuf().get(), downWidth, downHeight, shift);

  // upsample shifted
  Image::upsample22RGBA210(devUpBuffer.gpuBuf(), devDownBuffer.gpuBuf(), upWidth, upHeight, true,
                           GPU::Stream::getDefault());
  cudaStreamSynchronize(0);
  std::vector<uint32_t> hostShiftedUpBuffer;
  devUpBuffer.readback(hostShiftedUpBuffer);

  // shift the output twice as much as the input
  shiftHostLeft(hostUpBuffer.data(), upWidth, upHeight, 2 * shift);

  ENSURE_RGBA210_ARRAY_EQ(hostUpBuffer.data(), hostShiftedUpBuffer.data(), upWidth, upHeight);
}

void testUpsampleBench(const int64_t width, const int64_t height) {
  PackedDeviceBuffer inputBuffer(width, height);
  inputBuffer.fill(5, 10, 15);
  PackedDeviceBuffer outputBuffer(2 * width, 2 * height);
  auto potStream = GPU::Stream::create();
  {
    cudaDeviceSynchronize();
    Util::SimpleProfiler p("Upsample bench", false, Logger::get(Logger::Info));
    for (int i = 0; i < 300; ++i) {
      Image::upsample22RGBA210(outputBuffer.gpuBuf(), inputBuffer.gpuBuf(), 2 * width, 2 * height, false,
                               potStream.value());
    }
    cudaDeviceSynchronize();
  }
  potStream.value().destroy();
}

void testSubsampleBench(const int64_t width, const int64_t height) {
  PackedDeviceBuffer inputBuffer(width, height);
  inputBuffer.fill(5, 10, 15);
  PackedDeviceBuffer outputBuffer((width + 1) / 2, (height + 1) / 2);
  auto potStream = GPU::Stream::create();
  {
    cudaDeviceSynchronize();
    Util::SimpleProfiler p("Subsample bench", false, Logger::get(Logger::Info));
    for (int i = 0; i < 1000; ++i) {
      Image::subsample22RGBA(outputBuffer.gpuBuf(), inputBuffer.gpuBuf(), width, height, potStream.value());
    }
    cudaDeviceSynchronize();
  }
  potStream.value().destroy();
}
}  // namespace Testing
}  // namespace VideoStitch

#define PACKSAME(v, a) VideoStitch::Image::RGBA::pack((v), (v), (v), (a))
#define PACKSAME210(v, a) VideoStitch::Image::RGB210::pack((v), (v), (v), (a))
#define NA 0

int main(int argc, char** argv) {
  VideoStitch::Testing::initTest();
  VideoStitch::Testing::ENSURE(VideoStitch::GPU::setDefaultBackendDevice(0));

  // SubSample
  {
    unsigned char in[] = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16};
    unsigned char expOut[] = {
        3,
        5,
        11,
        13,
    };
    VideoStitch::Testing::testSubsample22(in, expOut, 4, 4);
  }
  {
    unsigned char in[] = {1, 2, 3, 4, 91, 5, 6, 7, 8, 92, 9, 10, 11, 12, 93, 13, 14, 15, 16, 94};
    unsigned char expOut[] = {3, 5, 91, 11, 13, 93};
    VideoStitch::Testing::testSubsample22(in, expOut, 5, 4);
  }
  {
    unsigned char in[] = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17, 18, 19, 20};
    unsigned char expOut[] = {3, 5, 11, 13, 17, 19};
    VideoStitch::Testing::testSubsample22(in, expOut, 4, 5);
  }
  {
    unsigned char in[] = {1, 2, 3, 4, 91, 5, 6, 7, 8, 92, 9, 10, 11, 12, 93, 13, 14, 15, 16, 94, 17, 18, 19, 20, 95};
    unsigned char expOut[] = {3, 5, 91, 11, 13, 93, 17, 19, 95};
    VideoStitch::Testing::testSubsample22(in, expOut, 5, 5);
  }

  /*
    //SubSample22Mask
    {
      uint32_t in[] = {
        1,  2,  3,  4, 91,
        5,  6,  7,  8, 92,
        9, 10, 11, 12, 93,
        13, 14, 15, 16, 94,
        17, 18, 19, 20, 95
      };
      uint32_t inMask[] = {
        1, 1, 1, 0, 0,
        1, 1, 1, 2, 2,
        1, 1, 1, 2, 2,
        3, 3, 3, 2, 2,
        3, 3, 3, 2, 2
      };
      uint32_t expOut[] = {
        3,  3, 91,
        9, 11, 93,
        17, 19, 95
      };
      uint32_t expOutMask[] = {
        1, 1, 0,
        1, 1, 2,
        3, 3, 2
      };
      // XXX TODO FIXME VideoStitch::Testing::testSubsample22Mask(in, inMask, expOut, expOutMask, 5, 5);
    }

    //SubSampleNearest
    {
      uint32_t in[] = {
        1,  2,  3,  4,
        5,  6,  7,  8,
        9, 10, 11, 12,
        13, 14, 15, 16
      };
      uint32_t expOut[] = {
        1,  3,
        9, 11,
      };
      // XXX TODO FIXME VideoStitch::Testing::testSubsample22Nearest(in, expOut, 4, 4);
    }
    {
      uint32_t in[] = {
        1,  2,  3,  4, 91,
        5,  6,  7,  8, 92,
        9, 10, 11, 12, 93,
        13, 14, 15, 16, 94
      };
      uint32_t expOut[] = {
        1,  3, 91,
        9, 11, 93
      };
      // XXX TODO FIXME VideoStitch::Testing::testSubsample22Nearest(in, expOut, 5, 4);
    }
    {
      uint32_t in[] = {
        1,  2,  3,  4,
        5,  6,  7,  8,
        9, 10, 11, 12,
        13, 14, 15, 16,
        17, 18, 19, 20
      };
      uint32_t expOut[] = {
        1,  3,
        9, 11,
        17, 19
      };
      // XXX TODO FIXME VideoStitch::Testing::testSubsample22Nearest(in, expOut, 4, 5);
    }
    {
      uint32_t in[] = {
        1,  2,  3,  4, 91,
        5,  6,  7,  8, 92,
        9, 10, 11, 12, 93,
        13, 14, 15, 16, 94,
        17, 18, 19, 20, 95
      };
      uint32_t expOut[] = {
        1,  3, 91,
        9, 11, 93,
        17, 19, 95
      };
      // XXX TODO FIXME VideoStitch::Testing::testSubsample22Nearest(in, expOut, 5, 5);
    }

    //SubSample mask
    {
      unsigned char in[] = {
        0, 0, 0, 0,
        0, 0, 0, 0,
        0, 0, 0, 0,
        0, 0, 0, 0
      };
      unsigned char expOut[] = {
        0, 0,
        0, 0,
      };
      // XXX TODO FIXME VideoStitch::Testing::testSubsampleMask22(in, expOut, 4, 4);
    }
    {
      unsigned char in[] = {
        1, 0, 0, 0,
        0, 0, 0, 0,
        0, 0, 0, 0,
        0, 0, 0, 0
      };
      unsigned char expOut[] = {
        1, 0,
        0, 0,
      };
      // XXX TODO FIXME VideoStitch::Testing::testSubsampleMask22(in, expOut, 4, 4);
    }
    {
      unsigned char in[] = {
        0, 0, 0, 1,
        0, 0, 0, 0,
        0, 0, 0, 0,
        0, 0, 0, 0
      };
      unsigned char expOut[] = {
        0, 1,
        0, 0,
      };
      // XXX TODO FIXME VideoStitch::Testing::testSubsampleMask22(in, expOut, 4, 4);
    }
    {
      unsigned char in[] = {
        0, 0, 0, 0,
        0, 0, 0, 0,
        0, 0, 0, 0,
        1, 0, 0, 0
      };
      unsigned char expOut[] = {
        0, 0,
        1, 0,
      };
      // XXX TODO FIXME VideoStitch::Testing::testSubsampleMask22(in, expOut, 4, 4);
    }
    {
      unsigned char in[] = {
        0, 0, 0, 0,
        0, 0, 0, 0,
        0, 0, 0, 0,
        0, 0, 0, 1
      };
      unsigned char expOut[] = {
        0, 0,
        0, 1,
      };
      // XXX TODO FIXME VideoStitch::Testing::testSubsampleMask22(in, expOut, 4, 4);
    }
    {
      unsigned char in[] = {
        0,  0,  0,  0,  0,
        0,  0,  0,  0,  0,
        0,  0,  0,  0,  1,
        0,  0,  0,  0,  0,
      };
      unsigned char expOut[] = {
        0, 0, 0,
        0, 0, 1
      };
      // XXX TODO FIXME VideoStitch::Testing::testSubsampleMask22(in, expOut, 5, 4);
    }
    {
      unsigned char in[] = {
        0,  0,  0,  0,  0,
        0,  0,  0,  0,  0,
        0,  0,  0,  1,  0,
        0,  0,  0,  0,  0,
      };
      unsigned char expOut[] = {
        0, 0, 0,
        0, 1, 0
      };
      // XXX TODO FIXME VideoStitch::Testing::testSubsampleMask22(in, expOut, 5, 4);
    }
  */

  // Subsample, with alpha
  uint32_t in[] = {PACKSAME(1, 0),    PACKSAME(2, 0),    PACKSAME(3, 0),    PACKSAME(4, 255),  PACKSAME(91, 0),
                   PACKSAME(5, 0),    PACKSAME(6, 0),    PACKSAME(7, 0),    PACKSAME(8, 0),    PACKSAME(92, 255),
                   PACKSAME(9, 0),    PACKSAME(10, 255), PACKSAME(11, 255), PACKSAME(12, 255), PACKSAME(93, 0),
                   PACKSAME(13, 255), PACKSAME(14, 0),   PACKSAME(15, 0),   PACKSAME(16, 255), PACKSAME(94, 0),
                   PACKSAME(17, 255), PACKSAME(18, 255), PACKSAME(19, 0),   PACKSAME(20, 0),   PACKSAME(95, 255)};
  uint32_t expOut[] = {PACKSAME(NA, 0), PACKSAME(4, 255),  PACKSAME(92, 255), PACKSAME(11, 255), PACKSAME(13, 255),
                       PACKSAME(NA, 0), PACKSAME(17, 255), PACKSAME(NA, 0),   PACKSAME(95, 255)};
  VideoStitch::Testing::testSubsample22RGBA(in, expOut, 5, 5);

  // UpSample
  {
    uint32_t expOut[] = {0, 4, 12, 16, 12, 16, 24, 28, 36, 40, 48, 52, 48, 52, 60, 64};
    uint32_t in[] = {
        0,
        16,
        48,
        64,
    };
    VideoStitch::Testing::testUpsample22(in, expOut, 4, 4);
  }
  {
    uint32_t in[] = {PACKSAME210(0, 1), PACKSAME210(16, 1), PACKSAME210(48, 1), PACKSAME210(64, 1)};
    uint32_t expOut[] = {PACKSAME210(0, 1),  PACKSAME210(4, 1),  PACKSAME210(12, 1), PACKSAME210(16, 1),
                         PACKSAME210(12, 1), PACKSAME210(16, 1), PACKSAME210(24, 1), PACKSAME210(28, 1),
                         PACKSAME210(36, 1), PACKSAME210(40, 1), PACKSAME210(48, 1), PACKSAME210(52, 1),
                         PACKSAME210(48, 1), PACKSAME210(52, 1), PACKSAME210(60, 1), PACKSAME210(64, 1)};
    VideoStitch::Testing::testUpsample22RGBA210(in, expOut, 4, 4);
  }
  {
    uint32_t in[] = {
        PACKSAME210(0, 1),  PACKSAME210(16, 1), PACKSAME210(32, 1),
        PACKSAME210(48, 1), PACKSAME210(64, 1), PACKSAME210(80, 1),
    };
    uint32_t expOut[] = {PACKSAME210(0, 1),  PACKSAME210(4, 1),  PACKSAME210(12, 1), PACKSAME210(20, 1),
                         PACKSAME210(28, 1), PACKSAME210(32, 1), PACKSAME210(12, 1), PACKSAME210(16, 1),
                         PACKSAME210(24, 1), PACKSAME210(32, 1), PACKSAME210(40, 1), PACKSAME210(44, 1),
                         PACKSAME210(36, 1), PACKSAME210(40, 1), PACKSAME210(48, 1), PACKSAME210(56, 1),
                         PACKSAME210(64, 1), PACKSAME210(68, 1), PACKSAME210(48, 1), PACKSAME210(52, 1),
                         PACKSAME210(60, 1), PACKSAME210(68, 1), PACKSAME210(76, 1), PACKSAME210(80, 1)};
    VideoStitch::Testing::testUpsample22RGBA210(in, expOut, 6, 4);
  }
  {
    uint32_t in[] = {PACKSAME210(0, 1),  PACKSAME210(48, 1), PACKSAME210(16, 1),
                     PACKSAME210(64, 1), PACKSAME210(32, 1), PACKSAME210(80, 1)};
    uint32_t expOut[] = {PACKSAME210(0, 1),  PACKSAME210(12, 1), PACKSAME210(36, 1), PACKSAME210(48, 1),
                         PACKSAME210(4, 1),  PACKSAME210(16, 1), PACKSAME210(40, 1), PACKSAME210(52, 1),
                         PACKSAME210(12, 1), PACKSAME210(24, 1), PACKSAME210(48, 1), PACKSAME210(60, 1),
                         PACKSAME210(20, 1), PACKSAME210(32, 1), PACKSAME210(56, 1), PACKSAME210(68, 1),
                         PACKSAME210(28, 1), PACKSAME210(40, 1), PACKSAME210(64, 1), PACKSAME210(76, 1),
                         PACKSAME210(32, 1), PACKSAME210(44, 1), PACKSAME210(68, 1), PACKSAME210(80, 1)};
    VideoStitch::Testing::testUpsample22RGBA210(in, expOut, 4, 6);
  }
  {
    uint32_t in[] = {
        PACKSAME210(0, 1),  PACKSAME210(16, 1), PACKSAME210(32, 1),
        PACKSAME210(48, 1), PACKSAME210(64, 1), PACKSAME210(80, 1),
    };
    uint32_t expOut[] = {
        PACKSAME210(0, 1),  PACKSAME210(4, 1),  PACKSAME210(12, 1), PACKSAME210(20, 1), PACKSAME210(28, 1),
        PACKSAME210(12, 1), PACKSAME210(16, 1), PACKSAME210(24, 1), PACKSAME210(32, 1), PACKSAME210(40, 1),
        PACKSAME210(36, 1), PACKSAME210(40, 1), PACKSAME210(48, 1), PACKSAME210(56, 1), PACKSAME210(64, 1),
        PACKSAME210(48, 1), PACKSAME210(52, 1), PACKSAME210(60, 1), PACKSAME210(68, 1), PACKSAME210(76, 1),
    };
    VideoStitch::Testing::testUpsample22RGBA210(in, expOut, 5, 4);
  }
  // upsample, with alpha
  {
    uint32_t in[] = {
        PACKSAME210(0, 0),  PACKSAME210(16, 0), PACKSAME210(32, 0),
        PACKSAME210(48, 1), PACKSAME210(64, 1), PACKSAME210(80, 1),
    };
    uint32_t expOut[] = {PACKSAME210(NA, 0), PACKSAME210(NA, 0), PACKSAME210(NA, 0), PACKSAME210(NA, 0),
                         PACKSAME210(NA, 0), PACKSAME210(48, 1), PACKSAME210(52, 1), PACKSAME210(60, 1),
                         PACKSAME210(68, 1), PACKSAME210(76, 1), PACKSAME210(48, 1), PACKSAME210(52, 1),
                         PACKSAME210(60, 1), PACKSAME210(68, 1), PACKSAME210(76, 1)};
    VideoStitch::Testing::testUpsample22RGBA210(in, expOut, 5, 3);
  }

  /*{
    unsigned w = 351;
    unsigned h = 432;
    unsigned char* data = new unsigned char[3 * w * h];
    int v = 7;
    for (unsigned i = 0; i < 3 * w * h; ++i) {
      data[i] = (unsigned char)v;
      v = (v * 13) % 255;
    }
    VideoStitch::Testing::testRoundTripsample(data, w, h, 16);
    delete[] data;
  }*/

  VideoStitch::Testing::testWrappingUpsample22RGBA210(2 * 16, 2 * 16, 16, 1);
  VideoStitch::Testing::testWrappingUpsample22RGBA210(2 * 16 + 2, 2 * 16, 16, 1);
  VideoStitch::Testing::testWrappingUpsample22RGBA210(2 * 16, 16 + 1, 2 * 16, 1);
  VideoStitch::Testing::testWrappingUpsample22RGBA210(2 * 16, 2 * 16 + 2, 16, 1);
  VideoStitch::Testing::testWrappingUpsample22RGBA210(2 * 16 + 2, 2 * 16 + 2, 16, 1);

  if (argc > 1 && !strcmp(argv[1], "bench_down")) {
    std::cout << "Benching subsampling, weird dims..." << std::endl;
    VideoStitch::Testing::testSubsampleBench(4215, 2111);
    std::cout << "Benching subsampling, power of two dims..." << std::endl;
    VideoStitch::Testing::testSubsampleBench(4096, 2048);
  }

  if (argc > 1 && !strcmp(argv[1], "bench_up")) {
    std::cout << "Benching upsampling, weird dims..." << std::endl;
    VideoStitch::Testing::testUpsampleBench(4215, 2111);
    std::cout << "Benching upsampling, power of two dims..." << std::endl;
    VideoStitch::Testing::testUpsampleBench(4096, 2048);
  }

  cudaDeviceReset();
  return 0;
}