1.临近点插值：

__global__ void nearest_BGR2RGB_nhwc2nchw_norm_kernel(
    float* tar, uint8_t* src, 
    int tarW, int tarH, 
    int srcW, int srcH,
    float scaled_w, float scaled_h,
    float* d_mean, float* d_std) 

    这次学习cuda，我发现一个技巧就是，在对矩阵或者图像进行计算时，我们的线程是面向结果的，就是说将线程安排好与结果相对，然后去反推计算过程，得到最初的数据与线程的关系。

有这个思想，就可以想到对于临近点插值，我们第一步就说把tar的二维索引与cuda的线程对齐。

int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

这两步得到thread的全局索引，我们把它也视为是tar的二维索引。

现在我们开始反推，去得到src_x与src_y.

无论是临近点插值还是双线性插值，我们第一步都是算scale_w和scale_h。

float scaled_h = (float)srcH / tarH;
float scaled_w = (float)srcW / tarW;

建议这一步可以放到最后，因为无论是临近还是双线性都会用到，没必要两次核函数都算一次。

有了scale后我们可以去还原坐标。

   int src_y = floor((float)y * scaled_h);
   int src_x = floor((float)x * scaled_w);

这里我们要注意的是，我们求出来的这两个坐标都是比真实的小的。由于floor函数直接把浮点变成整型了，他会直接无视小数。

此时两张图的x，y已经有了映射关系，就差遍历了。

这次我们不仅做了resize，而且在遍历的时候要把BGR2RGB，nhwc2nchw，norm这三个算子也融合在一起。

这里面要讲一件事情，对于opencv来说一张图片的格式是[rgb,rgb,rgb......],而我们在tensorrt还是别的架构，一般是[rrrrrrrrr...,gggggggggg.......,bbbbbbbb.......],所以两个的遍历方式也是不同的，我们要做的是要把src的r，g，b都单独挑出来，然后齐齐整整的放到tar中。

对于tar来说，第一步，对于每张图片，我们都要把二维的坐标去转换为一维的索引。这里有个公式：

int tarIdx  = y * tarW + x;

这个公式把我们一张二维图片的所有坐标转化为索引。

我们要分成3个空间，每个空间的大小为一张图片的size。也就是说一张图片的size应该是

（h*w*c*sizeof（float））才能储存一张图片数据。

那仔细想想我们要把tar分成3段，每段大小为[h*w].则：

 int tarArea = tarW * tarH;

现在我们如果存储r的像素值：则是

tar[tar_idx+tar_area*0]

如果是g则把0换成1，b则是换成2.

现在对于src来说，我们要去拿:由于结构不同，我们遍历的方式也不同。

但是第一步都是去把二维转化为一维：

src_idx = src_y*src_w +src_x

如果我们想拿g，bgrbgrbgr.....,这三个循环，第一个g的索引为1，第二个为4，第三个为7，发现规律就是：1+src_idx*3

如果是r能，则是2+src_idx*3,发现就说最初索引值+一维索引值×3，我们可以把一维索引值*3看成一步走多宽，如果*1，则走的很短，而×3则恰好每一步都走到相同的像素上。

两个遍历我们都完成了，剩下就说直接去赋值了

     tar[tarIdx + tarArea * 0] = (src[srcIdx + 2] / 255.0f - d_mean[2]) / d_std[2];
     tar[tarIdx + tarArea * 1] = (src[srcIdx + 1] / 255.0f - d_mean[1]) / d_std[1];
     tar[tarIdx + tarArea * 2] = (src[srcIdx + 0] / 255.0f - d_mean[0]) / d_std[0];

这一步我们不仅做了赋值，也把BGR2RGB，nhwc2nchw，norm给做了。

__global__ void nearest_BGR2RGB_nhwc2nchw_norm_kernel(
    float* tar, uint8_t* src, 
    int tarW, int tarH, 
    int srcW, int srcH,
    float scaled_w, float scaled_h,
    float* d_mean, float* d_std) 
{
    // nearest neighbour -- resized之后的图tar上的坐标
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;

    // nearest neighbour -- 计算最近坐标
    int src_y = floor((float)y * scaled_h);
    int src_x = floor((float)x * scaled_w);

    if (src_x < 0 || src_y < 0 || src_x > srcW || src_y > srcH) {
        // nearest neighbour -- 对于越界的部分，不进行计算
    } else {
        // nearest neighbour -- 计算tar中对应坐标的索引
        int tarIdx  = y * tarW + x;
        int tarArea = tarW * tarH;

        // nearest neighbour -- 计算src中最近邻坐标的索引
        int srcIdx = (src_y * srcW + src_x) * 3;

        // nearest neighbour -- 实现nearest beighbour的resize + BGR2RGB + nhwc2nchw + norm
        tar[tarIdx + tarArea * 0] = (src[srcIdx + 2] / 255.0f - d_mean[2]) / d_std[2];
        tar[tarIdx + tarArea * 1] = (src[srcIdx + 1] / 255.0f - d_mean[1]) / d_std[1];
        tar[tarIdx + tarArea * 2] = (src[srcIdx + 0] / 255.0f - d_mean[0]) / d_std[0];
    }
}

现在来做双线性插值，我们发现直接找临近值明显误差很大，不平滑，而双线性插值则解决了这个问题：

int src_y1 = floor((y + 0.5) * scaled_h - 0.5);
    int src_x1 = floor((x + 0.5) * scaled_w - 0.5);
    int src_y2 = src_y1 + 1;
    int src_x2 = src_x1 + 1;

我们不仅优化了最初的临近点的选择，又增加了一个坐标。根据排列组合，现在有4个点，那该如何处理把这个像素值完美的赋值给tar呢，我们去算一个权重，这个权重能算这4个点到底那个与tar（x，y）有关。原理是这样的

我们看阴影的面积，阴影面积越大代表什么，p点离a点越近，所以点对应的面积就是这个权重，我们可以算出来ABCD这四个点对应的面积也就是权重。

        // bilinear interpolation -- 计算原图上的坐标(浮点类型)在0~1之间的值
        float th   = ((y + 0.5) * scaled_h - 0.5) - src_y1;
        float tw   = ((x + 0.5) * scaled_w - 0.5) - src_x1;

        // bilinear interpolation -- 计算面积(这里建议自己手画一张图来理解一下)
        float a1_1 = (1.0 - tw) * (1.0 - th);  //右下
        float a1_2 = tw * (1.0 - th);          //左下
        float a2_1 = (1.0 - tw) * th;          //右上
        float a2_2 = tw * th;                  //左上

我们算出来每个点对tar（x，y）的重要性了，现在对于tar来说遍历还是那样遍历

    int tarIdx    = y * tarW  + x;
    int tarArea   = tarW * tarH;

但是我们要把ABCD这4个点的二维都转成一维的索引，这四个像素值×权重后加一起得到的那个最终的像素值就是最适合tar（x，y）的值：

这是4个点对应的一维索引

 int srcIdx1_1 = (src_y1 * srcW + src_x1) * 3;  //左上
        int srcIdx1_2 = (src_y1 * srcW + src_x2) * 3;  //右上
        int srcIdx2_1 = (src_y2 * srcW + src_x1) * 3;  //左下
        int srcIdx2_2 = (src_y2 * srcW + src_x2) * 3;  //右下

现在我们索引结束了，剩下的就是传值了

 
        tar[tarIdx + tarArea * 0] = 
            (round((a1_1 * src[srcIdx1_1 + 2] + 
                   a1_2 * src[srcIdx1_2 + 2] +
                   a2_1 * src[srcIdx2_1 + 2] +
                   a2_2 * src[srcIdx2_2 + 2])) / 255.0f - d_mean[2]) / d_std[2];

        tar[tarIdx + tarArea * 1] = 
            (round((a1_1 * src[srcIdx1_1 + 1] + 
                   a1_2 * src[srcIdx1_2 + 1] +
                   a2_1 * src[srcIdx2_1 + 1] +
                   a2_2 * src[srcIdx2_2 + 1])) / 255.0f - d_mean[1]) / d_std[1];

        tar[tarIdx + tarArea * 2] = 
            (round((a1_1 * src[srcIdx1_1 + 0] + 
                   a1_2 * src[srcIdx1_2 + 0] +
                   a2_1 * src[srcIdx2_1 + 0] +
                   a2_2 * src[srcIdx2_2 + 0])) / 255.0f - d_mean[0]) / d_std[0];

似乎不需要用round取证会更好。