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取证会更好。
