#include "gpu_preprocess.hpp"
#include <algorithm>
#include <cmath>
namespace cms {
__global__ void yoloNormalize(float* out,unsigned char const* in,int pixels) {
    int i=blockDim.x*blockIdx.x+threadIdx.x;if(i>=pixels)return;
    out[i]=__fdiv_rn(float(in[3*i+2]),255.F);out[pixels+i]=__fdiv_rn(float(in[3*i+1]),255.F);out[2*pixels+i]=__fdiv_rn(float(in[3*i]),255.F);
}
void YoloPreprocess::enqueue(cv::Mat const& frame,float* destination,cudaStream_t stream) {
    double scale=std::min(double(width_)/frame.cols,double(height_)/frame.rows);
    int rw=int(std::nearbyint(frame.cols*scale)),rh=int(std::nearbyint(frame.rows*scale));
    int left=int(std::nearbyint((width_-rw)/2.0-.1)),top=int(std::nearbyint((height_-rh)/2.0-.1));
        size_t bytes=size_t(frame.cols)*frame.rows*3;if(capacity_<bytes){if(source_)CUDA_CHECK(cudaFree(source_));CUDA_CHECK(cudaMalloc(&source_,bytes));capacity_=bytes;}
        CUDA_CHECK(cudaMemcpy2DAsync(source_,frame.cols*3,frame.data,frame.step,frame.cols*3,frame.rows,cudaMemcpyHostToDevice,stream));
        auto c=nppContext(stream);Npp8u fill[3]={114,114,114};nppCheck(nppiSet_8u_C3R_Ctx(fill,letter_,width_*3,{width_,height_},c));
        nppCheck(nppiResizeSqrPixel_8u_C3R_Ctx(source_,{frame.cols,frame.rows},frame.cols*3,{0,0,frame.cols,frame.rows},letter_,width_*3,{left,top,rw,rh},double(rw)/frame.cols,double(rh)/frame.rows,left,top,NPPI_INTER_LINEAR,c));
    int pixels=width_*height_;yoloNormalize<<<(pixels+255)/256,256,0,stream>>>(destination,letter_,pixels);CUDA_CHECK(cudaGetLastError());
}
}