[cpp] elena cuda

Viewer

copydownloadembedprintName: elena cuda
  1.  
  2. #include <iostream>
  3. #include <cstdio>
  4. #include <cstdlib>
  5. #include <cstring>
  6. #define EQUAL(a,b) (strcmp((a),(b))==0)
  7. #define ABORT(Msg)                                                   \
  8.   {                                                                        \
  9.     std::cerr << ": \033[1;91m"                                            \
  10.               << "[Fatal]"                                                 \
  11.               << "\033[m " << __FILE__ << ": " << __FUNCTION__ << ": Line" \
  12.               << __LINE__ << ": " << Msg << std::endl;                     \
  13.     std::abort();                                                               \
  14.   }
  15.  
  16. #include "elena_int.h"
  17.  
  18.  
  19. #include <cuda_runtime.h>
  20. #define cuErrCheck(res)                                        \
  21.     {                                                          \
  22.         if (res != cudaSuccess)                                \
  23.             ABORT("cuda assert: " << cudaGetErrorString(res)); \
  24.     }
  25. #define BLOCK_SIZE 128
  26.  
  27. #include "elena_registry.h"
  28. namespace cuda_fa993066c1ec32a14d8d06ecfb8d5c54e7a04ec31ba833e1abf1f16c05e027c7 {
  29.  __global__ void BGR_Bilinear_Kernel(uint64_t resize_h, uint64_t resize_w, uint64_t crop_h, uint64_t crop_w, int32_t crop_top, int32_t crop_left, float norm_mean_0, float norm_mean_1, float norm_mean_2, float norm_std_0, float norm_std_1, float norm_std_2, uint64_t pad_h, uint64_t pad_w, int32_t pad_top, int32_t pad_left, int32_t pad_bottom, int32_t pad_right, float pad_value, int16_t* __restrict__ cubfh, int16_t* __restrict__ cubfw, int32_t* __restrict__ inth, int32_t* __restrict__ intw, uint8_t* __restrict__ input, float* __restrict__ HWC2CHW, uint64_t h, uint64_t w) {
  30.     float BGR_std[3];
  31.     float BGR_mean[3];
  32.     for (uint64_t iter3 = 0; iter3 < 0 + 3; ++iter3) {
  33.         if ((iter3 < ((0) + (3)))) {
  34.             BGR_std[((iter3) - (0))] = ((iter3 == 0) ? norm_std_0 : ((iter3 == 1) ? norm_std_1 : norm_std_2));
  35.         }
  36.     }
  37.     for (uint64_t iter2 = 0; iter2 < 0 + 3; ++iter2) {
  38.         if ((iter2 < ((0) + (3)))) {
  39.             BGR_mean[((iter2) - (0))] = ((iter2 == 0) ? norm_mean_0 : ((iter2 == 1) ? norm_mean_1 : norm_mean_2));
  40.         }
  41.     }
  42.     const uint64_t iter19_iter20_fused_iter21_fused = ((((((blockIdx.x) * (128))) + (threadIdx.x))) + (0));
  43.     const uint64_t iter21 = (((iter19_iter20_fused_iter21_fused % 3)) + (0));
  44.     const uint64_t iter19_iter20_fused = ((((iter19_iter20_fused_iter21_fused) / (3))) + (0));
  45.     const uint64_t iter20 = (((iter19_iter20_fused % crop_w)) + (0));
  46.     const uint64_t iter19 = ((((iter19_iter20_fused) / (crop_w))) + (0));
  47.     if ((iter19 < ((0) + (crop_h)))) {
  48.         if ((iter20 < ((0) + (crop_w)))) {
  49.             if ((iter21 < ((0) + (3)))) {
  50.                 HWC2CHW[((((((iter20) - (0))) + (((crop_w) * (((iter19) - (0))))))) + (((((crop_w) * (crop_h))) * (((iter21) - (0))))))] = (((((float)(((((((((((((((cubfh[((((iter19) + (crop_top))) + (((resize_h) * (0))))]) * (cubfw[((((iter20) + (crop_left))) + (((resize_w) * (0))))]))) * (input[((((((2) - (iter21))) + (((3) * (intw[((((iter20) + (crop_left))) + (((resize_w) * (0))))]))))) + (((((3) * (w))) * (inth[((((iter19) + (crop_top))) + (((resize_h) * (0))))]))))]))) + (((((cubfh[((((iter19) + (crop_top))) + (((resize_h) * (1))))]) * (cubfw[((((iter20) + (crop_left))) + (((resize_w) * (0))))]))) * (input[((((((2) - (iter21))) + (((3) * (intw[((((iter20) + (crop_left))) + (((resize_w) * (0))))]))))) + (((((3) * (w))) * (inth[((((iter19) + (crop_top))) + (((resize_h) * (1))))]))))]))))) + (((((cubfh[((((iter19) + (crop_top))) + (((resize_h) * (0))))]) * (cubfw[((((iter20) + (crop_left))) + (((resize_w) * (1))))]))) * (input[((((((2) - (iter21))) + (((3) * (intw[((((iter20) + (crop_left))) + (((resize_w) * (1))))]))))) + (((((3) * (w))) * (inth[((((iter19) + (crop_top))) + (((resize_h) * (0))))]))))]))))) + (((((cubfh[((((iter19) + (crop_top))) + (((resize_h) * (1))))]) * (cubfw[((((iter20) + (crop_left))) + (((resize_w) * (1))))]))) * (input[((((((2) - (iter21))) + (((3) * (intw[((((iter20) + (crop_left))) + (((resize_w) * (1))))]))))) + (((((3) * (w))) * (inth[((((iter19) + (crop_top))) + (((resize_h) * (1))))]))))]))))) + (2097152))) / (4194304)))) - (BGR_mean[((iter21) - (0))]))) / (BGR_std[((iter21) - (0))]));
  51.             }
  52.         }
  53.     }
  54. }
  55.  
  56.  __global__ void BGR_Nearest_Kernel(uint64_t resize_h, uint64_t resize_w, uint64_t crop_h, uint64_t crop_w, int32_t crop_top, int32_t crop_left, float norm_mean_0, float norm_mean_1, float norm_mean_2, float norm_std_0, float norm_std_1, float norm_std_2, uint64_t pad_h, uint64_t pad_w, int32_t pad_top, int32_t pad_left, int32_t pad_bottom, int32_t pad_right, float pad_value, uint8_t* __restrict__ input, float* __restrict__ HWC2CHW, uint64_t h, uint64_t w) {
  57.     float BGR_std[3];
  58.     float scale[2];
  59.     float BGR_mean[3];
  60.     for (uint64_t iter28 = 0; iter28 < 0 + 3; ++iter28) {
  61.         if ((iter28 < ((0) + (3)))) {
  62.             BGR_std[((iter28) - (0))] = ((iter28 == 0) ? norm_std_0 : ((iter28 == 1) ? norm_std_1 : norm_std_2));
  63.         }
  64.     }
  65.     for (uint64_t iter32 = 0; iter32 < 0 + 2; ++iter32) {
  66.         if ((iter32 < ((0) + (2)))) {
  67.             scale[((iter32) - (0))] = ((iter32 == 0) ? (((float)(h)) / ((float)(resize_h))) : (((float)(w)) / ((float)(resize_w))));
  68.         }
  69.     }
  70.     for (uint64_t iter27 = 0; iter27 < 0 + 3; ++iter27) {
  71.         if ((iter27 < ((0) + (3)))) {
  72.             BGR_mean[((iter27) - (0))] = ((iter27 == 0) ? norm_mean_0 : ((iter27 == 1) ? norm_mean_1 : norm_mean_2));
  73.         }
  74.     }
  75.     const uint64_t iter45_iter46_fused_iter47_fused = ((((((blockIdx.x) * (128))) + (threadIdx.x))) + (0));
  76.     const uint64_t iter47 = (((iter45_iter46_fused_iter47_fused % 3)) + (0));
  77.     const uint64_t iter45_iter46_fused = ((((iter45_iter46_fused_iter47_fused) / (3))) + (0));
  78.     const uint64_t iter46 = (((iter45_iter46_fused % crop_w)) + (0));
  79.     const uint64_t iter45 = ((((iter45_iter46_fused) / (crop_w))) + (0));
  80.     if ((iter45 < ((0) + (crop_h)))) {
  81.         if ((iter46 < ((0) + (crop_w)))) {
  82.             if ((iter47 < ((0) + (3)))) {
  83.                 HWC2CHW[((((((iter46) - (0))) + (((crop_w) * (((iter45) - (0))))))) + (((((crop_w) * (crop_h))) * (((iter47) - (0))))))] = (((((float)(input[((((((2) - (iter47))) + (((3) * (min((uint64_t)(static_cast<uint64_t>(floorf(((((iter46) + (crop_left))) * (scale[((1) - (0))]))))), (uint64_t)((w) - (1)))))))) + (((((3) * (w))) * (min((uint64_t)(static_cast<uint64_t>(floorf(((((iter45) + (crop_top))) * (scale[((0) - (0))]))))), (uint64_t)((h) - (1)))))))])) - (BGR_mean[((iter47) - (0))]))) / (BGR_std[((iter47) - (0))]));
  84.             }
  85.         }
  86.     }
  87. }
  88.  
  89.  __global__ void RGB_Bilinear_Kernel(uint64_t resize_h, uint64_t resize_w, uint64_t crop_h, uint64_t crop_w, int32_t crop_top, int32_t crop_left, float norm_mean_0, float norm_mean_1, float norm_mean_2, float norm_std_0, float norm_std_1, float norm_std_2, uint64_t pad_h, uint64_t pad_w, int32_t pad_top, int32_t pad_left, int32_t pad_bottom, int32_t pad_right, float pad_value, int16_t* __restrict__ cubfh, int16_t* __restrict__ cubfw, int32_t* __restrict__ inth, int32_t* __restrict__ intw, uint8_t* __restrict__ input, float* __restrict__ HWC2CHW, uint64_t h, uint64_t w) {
  90.     float BGR_std[3];
  91.     float BGR_mean[3];
  92.     for (uint64_t iter54 = 0; iter54 < 0 + 3; ++iter54) {
  93.         if ((iter54 < ((0) + (3)))) {
  94.             BGR_std[((iter54) - (0))] = ((iter54 == 0) ? norm_std_0 : ((iter54 == 1) ? norm_std_1 : norm_std_2));
  95.         }
  96.     }
  97.     for (uint64_t iter53 = 0; iter53 < 0 + 3; ++iter53) {
  98.         if ((iter53 < ((0) + (3)))) {
  99.             BGR_mean[((iter53) - (0))] = ((iter53 == 0) ? norm_mean_0 : ((iter53 == 1) ? norm_mean_1 : norm_mean_2));
  100.         }
  101.     }
  102.     const uint64_t iter73_iter74_fused_iter75_fused = ((((((blockIdx.x) * (128))) + (threadIdx.x))) + (0));
  103.     const uint64_t iter75 = (((iter73_iter74_fused_iter75_fused % 3)) + (0));
  104.     const uint64_t iter73_iter74_fused = ((((iter73_iter74_fused_iter75_fused) / (3))) + (0));
  105.     const uint64_t iter74 = (((iter73_iter74_fused % crop_w)) + (0));
  106.     const uint64_t iter73 = ((((iter73_iter74_fused) / (crop_w))) + (0));
  107.     if ((iter73 < ((0) + (crop_h)))) {
  108.         if ((iter74 < ((0) + (crop_w)))) {
  109.             if ((iter75 < ((0) + (3)))) {
  110.                 HWC2CHW[((((((iter74) - (0))) + (((crop_w) * (((iter73) - (0))))))) + (((((crop_w) * (crop_h))) * (((iter75) - (0))))))] = (((((float)(((((((((((((((cubfh[((((iter73) + (crop_top))) + (((resize_h) * (0))))]) * (cubfw[((((iter74) + (crop_left))) + (((resize_w) * (0))))]))) * (input[((((((2) - (((2) - (iter75))))) + (((3) * (intw[((((iter74) + (crop_left))) + (((resize_w) * (0))))]))))) + (((((3) * (w))) * (inth[((((iter73) + (crop_top))) + (((resize_h) * (0))))]))))]))) + (((((cubfh[((((iter73) + (crop_top))) + (((resize_h) * (1))))]) * (cubfw[((((iter74) + (crop_left))) + (((resize_w) * (0))))]))) * (input[((((((2) - (((2) - (iter75))))) + (((3) * (intw[((((iter74) + (crop_left))) + (((resize_w) * (0))))]))))) + (((((3) * (w))) * (inth[((((iter73) + (crop_top))) + (((resize_h) * (1))))]))))]))))) + (((((cubfh[((((iter73) + (crop_top))) + (((resize_h) * (0))))]) * (cubfw[((((iter74) + (crop_left))) + (((resize_w) * (1))))]))) * (input[((((((2) - (((2) - (iter75))))) + (((3) * (intw[((((iter74) + (crop_left))) + (((resize_w) * (1))))]))))) + (((((3) * (w))) * (inth[((((iter73) + (crop_top))) + (((resize_h) * (0))))]))))]))))) + (((((cubfh[((((iter73) + (crop_top))) + (((resize_h) * (1))))]) * (cubfw[((((iter74) + (crop_left))) + (((resize_w) * (1))))]))) * (input[((((((2) - (((2) - (iter75))))) + (((3) * (intw[((((iter74) + (crop_left))) + (((resize_w) * (1))))]))))) + (((((3) * (w))) * (inth[((((iter73) + (crop_top))) + (((resize_h) * (1))))]))))]))))) + (2097152))) / (4194304)))) - (BGR_mean[((iter75) - (0))]))) / (BGR_std[((iter75) - (0))]));
  111.             }
  112.         }
  113.     }
  114. }
  115.  
  116.  __global__ void RGB_Nearest_Kernel(uint64_t resize_h, uint64_t resize_w, uint64_t crop_h, uint64_t crop_w, int32_t crop_top, int32_t crop_left, float norm_mean_0, float norm_mean_1, float norm_mean_2, float norm_std_0, float norm_std_1, float norm_std_2, uint64_t pad_h, uint64_t pad_w, int32_t pad_top, int32_t pad_left, int32_t pad_bottom, int32_t pad_right, float pad_value, uint8_t* __restrict__ input, float* __restrict__ HWC2CHW, uint64_t h, uint64_t w) {
  117.     float BGR_std[3];
  118.     float scale[2];
  119.     float BGR_mean[3];
  120.     for (uint64_t iter82 = 0; iter82 < 0 + 3; ++iter82) {
  121.         if ((iter82 < ((0) + (3)))) {
  122.             BGR_std[((iter82) - (0))] = ((iter82 == 0) ? norm_std_0 : ((iter82 == 1) ? norm_std_1 : norm_std_2));
  123.         }
  124.     }
  125.     for (uint64_t iter89 = 0; iter89 < 0 + 2; ++iter89) {
  126.         if ((iter89 < ((0) + (2)))) {
  127.             scale[((iter89) - (0))] = ((iter89 == 0) ? (((float)(h)) / ((float)(resize_h))) : (((float)(w)) / ((float)(resize_w))));
  128.         }
  129.     }
  130.     for (uint64_t iter81 = 0; iter81 < 0 + 3; ++iter81) {
  131.         if ((iter81 < ((0) + (3)))) {
  132.             BGR_mean[((iter81) - (0))] = ((iter81 == 0) ? norm_mean_0 : ((iter81 == 1) ? norm_mean_1 : norm_mean_2));
  133.         }
  134.     }
  135.     const uint64_t iter102_iter103_fused_iter104_fused = ((((((blockIdx.x) * (128))) + (threadIdx.x))) + (0));
  136.     const uint64_t iter104 = (((iter102_iter103_fused_iter104_fused % 3)) + (0));
  137.     const uint64_t iter102_iter103_fused = ((((iter102_iter103_fused_iter104_fused) / (3))) + (0));
  138.     const uint64_t iter103 = (((iter102_iter103_fused % crop_w)) + (0));
  139.     const uint64_t iter102 = ((((iter102_iter103_fused) / (crop_w))) + (0));
  140.     if ((iter102 < ((0) + (crop_h)))) {
  141.         if ((iter103 < ((0) + (crop_w)))) {
  142.             if ((iter104 < ((0) + (3)))) {
  143.                 HWC2CHW[((((((iter103) - (0))) + (((crop_w) * (((iter102) - (0))))))) + (((((crop_w) * (crop_h))) * (((iter104) - (0))))))] = (((((float)(input[((((((2) - (((2) - (iter104))))) + (((3) * (min((uint64_t)(static_cast<uint64_t>(floorf(((((iter103) + (crop_left))) * (scale[((1) - (0))]))))), (uint64_t)((w) - (1)))))))) + (((((3) * (w))) * (min((uint64_t)(static_cast<uint64_t>(floorf(((((iter102) + (crop_top))) * (scale[((0) - (0))]))))), (uint64_t)((h) - (1)))))))])) - (BGR_mean[((iter104) - (0))]))) / (BGR_std[((iter104) - (0))]));
  144.             }
  145.         }
  146.     }
  147. }
  148.  
  149.  __global__ void GRAY_Bilinear_Kernel(uint64_t resize_h, uint64_t resize_w, uint64_t crop_h, uint64_t crop_w, int32_t crop_top, int32_t crop_left, float norm_mean_0, float norm_mean_1, float norm_mean_2, float norm_std_0, float norm_std_1, float norm_std_2, uint64_t pad_h, uint64_t pad_w, int32_t pad_top, int32_t pad_left, int32_t pad_bottom, int32_t pad_right, float pad_value, int16_t* __restrict__ cubfh, int16_t* __restrict__ cubfw, int32_t* __restrict__ inth, int32_t* __restrict__ intw, uint8_t* __restrict__ input, float* __restrict__ HWC2CHW, uint64_t h, uint64_t w) {
  150.     float BGR_std[3];
  151.     float BGR_mean[3];
  152.     for (uint64_t iter111 = 0; iter111 < 0 + 3; ++iter111) {
  153.         if ((iter111 < ((0) + (3)))) {
  154.             BGR_std[((iter111) - (0))] = ((iter111 == 0) ? norm_std_0 : ((iter111 == 1) ? norm_std_1 : norm_std_2));
  155.         }
  156.     }
  157.     for (uint64_t iter110 = 0; iter110 < 0 + 3; ++iter110) {
  158.         if ((iter110 < ((0) + (3)))) {
  159.             BGR_mean[((iter110) - (0))] = ((iter110 == 0) ? norm_mean_0 : ((iter110 == 1) ? norm_mean_1 : norm_mean_2));
  160.         }
  161.     }
  162.     const uint64_t iter130_iter131_fused_iter132_fused = ((((((blockIdx.x) * (128))) + (threadIdx.x))) + (0));
  163.     const uint64_t iter132 = (((iter130_iter131_fused_iter132_fused % 3)) + (0));
  164.     const uint64_t iter130_iter131_fused = ((((iter130_iter131_fused_iter132_fused) / (3))) + (0));
  165.     const uint64_t iter131 = (((iter130_iter131_fused % crop_w)) + (0));
  166.     const uint64_t iter130 = ((((iter130_iter131_fused) / (crop_w))) + (0));
  167.     if ((iter130 < ((0) + (crop_h)))) {
  168.         if ((iter131 < ((0) + (crop_w)))) {
  169.             if ((iter132 < ((0) + (3)))) {
  170.                 HWC2CHW[((((((iter131) - (0))) + (((crop_w) * (((iter130) - (0))))))) + (((((crop_w) * (crop_h))) * (((iter132) - (0))))))] = (((((float)(((((((((((((((cubfh[((((iter130) + (crop_top))) + (((resize_h) * (0))))]) * (cubfw[((((iter131) + (crop_left))) + (((resize_w) * (0))))]))) * (input[((((0) + (((1) * (intw[((((iter131) + (crop_left))) + (((resize_w) * (0))))]))))) + (((((1) * (w))) * (inth[((((iter130) + (crop_top))) + (((resize_h) * (0))))]))))]))) + (((((cubfh[((((iter130) + (crop_top))) + (((resize_h) * (1))))]) * (cubfw[((((iter131) + (crop_left))) + (((resize_w) * (0))))]))) * (input[((((0) + (((1) * (intw[((((iter131) + (crop_left))) + (((resize_w) * (0))))]))))) + (((((1) * (w))) * (inth[((((iter130) + (crop_top))) + (((resize_h) * (1))))]))))]))))) + (((((cubfh[((((iter130) + (crop_top))) + (((resize_h) * (0))))]) * (cubfw[((((iter131) + (crop_left))) + (((resize_w) * (1))))]))) * (input[((((0) + (((1) * (intw[((((iter131) + (crop_left))) + (((resize_w) * (1))))]))))) + (((((1) * (w))) * (inth[((((iter130) + (crop_top))) + (((resize_h) * (0))))]))))]))))) + (((((cubfh[((((iter130) + (crop_top))) + (((resize_h) * (1))))]) * (cubfw[((((iter131) + (crop_left))) + (((resize_w) * (1))))]))) * (input[((((0) + (((1) * (intw[((((iter131) + (crop_left))) + (((resize_w) * (1))))]))))) + (((((1) * (w))) * (inth[((((iter130) + (crop_top))) + (((resize_h) * (1))))]))))]))))) + (2097152))) / (4194304)))) - (BGR_mean[((iter132) - (0))]))) / (BGR_std[((iter132) - (0))]));
  171.             }
  172.         }
  173.     }
  174. }
  175.  
  176.  __global__ void GRAY_Nearest_Kernel(uint64_t resize_h, uint64_t resize_w, uint64_t crop_h, uint64_t crop_w, int32_t crop_top, int32_t crop_left, float norm_mean_0, float norm_mean_1, float norm_mean_2, float norm_std_0, float norm_std_1, float norm_std_2, uint64_t pad_h, uint64_t pad_w, int32_t pad_top, int32_t pad_left, int32_t pad_bottom, int32_t pad_right, float pad_value, uint8_t* __restrict__ input, float* __restrict__ HWC2CHW, uint64_t h, uint64_t w) {
  177.     float BGR_std[3];
  178.     float scale[2];
  179.     float BGR_mean[3];
  180.     for (uint64_t iter139 = 0; iter139 < 0 + 3; ++iter139) {
  181.         if ((iter139 < ((0) + (3)))) {
  182.             BGR_std[((iter139) - (0))] = ((iter139 == 0) ? norm_std_0 : ((iter139 == 1) ? norm_std_1 : norm_std_2));
  183.         }
  184.     }
  185.     for (uint64_t iter146 = 0; iter146 < 0 + 2; ++iter146) {
  186.         if ((iter146 < ((0) + (2)))) {
  187.             scale[((iter146) - (0))] = ((iter146 == 0) ? (((float)(h)) / ((float)(resize_h))) : (((float)(w)) / ((float)(resize_w))));
  188.         }
  189.     }
  190.     for (uint64_t iter138 = 0; iter138 < 0 + 3; ++iter138) {
  191.         if ((iter138 < ((0) + (3)))) {
  192.             BGR_mean[((iter138) - (0))] = ((iter138 == 0) ? norm_mean_0 : ((iter138 == 1) ? norm_mean_1 : norm_mean_2));
  193.         }
  194.     }
  195.     const uint64_t iter159_iter160_fused_iter161_fused = ((((((blockIdx.x) * (128))) + (threadIdx.x))) + (0));
  196.     const uint64_t iter161 = (((iter159_iter160_fused_iter161_fused % 3)) + (0));
  197.     const uint64_t iter159_iter160_fused = ((((iter159_iter160_fused_iter161_fused) / (3))) + (0));
  198.     const uint64_t iter160 = (((iter159_iter160_fused % crop_w)) + (0));
  199.     const uint64_t iter159 = ((((iter159_iter160_fused) / (crop_w))) + (0));
  200.     if ((iter159 < ((0) + (crop_h)))) {
  201.         if ((iter160 < ((0) + (crop_w)))) {
  202.             if ((iter161 < ((0) + (3)))) {
  203.                 HWC2CHW[((((((iter160) - (0))) + (((crop_w) * (((iter159) - (0))))))) + (((((crop_w) * (crop_h))) * (((iter161) - (0))))))] = (((((float)(input[((((0) + (((1) * (min((uint64_t)(static_cast<uint64_t>(floorf(((((iter160) + (crop_left))) * (scale[((1) - (0))]))))), (uint64_t)((w) - (1)))))))) + (((((1) * (w))) * (min((uint64_t)(static_cast<uint64_t>(floorf(((((iter159) + (crop_top))) * (scale[((0) - (0))]))))), (uint64_t)((h) - (1)))))))])) - (BGR_mean[((iter161) - (0))]))) / (BGR_std[((iter161) - (0))]));
  204.             }
  205.         }
  206.     }
  207. }
  208.  
  209.  __global__ void BGRA_Bilinear_Kernel(uint64_t resize_h, uint64_t resize_w, uint64_t crop_h, uint64_t crop_w, int32_t crop_top, int32_t crop_left, float norm_mean_0, float norm_mean_1, float norm_mean_2, float norm_std_0, float norm_std_1, float norm_std_2, uint64_t pad_h, uint64_t pad_w, int32_t pad_top, int32_t pad_left, int32_t pad_bottom, int32_t pad_right, float pad_value, int16_t* __restrict__ cubfh, int16_t* __restrict__ cubfw, int32_t* __restrict__ inth, int32_t* __restrict__ intw, uint8_t* __restrict__ input, float* __restrict__ HWC2CHW, uint64_t h, uint64_t w) {
  210.     float BGR_std[3];
  211.     float BGR_mean[3];
  212.     for (uint64_t iter168 = 0; iter168 < 0 + 3; ++iter168) {
  213.         if ((iter168 < ((0) + (3)))) {
  214.             BGR_std[((iter168) - (0))] = ((iter168 == 0) ? norm_std_0 : ((iter168 == 1) ? norm_std_1 : norm_std_2));
  215.         }
  216.     }
  217.     for (uint64_t iter167 = 0; iter167 < 0 + 3; ++iter167) {
  218.         if ((iter167 < ((0) + (3)))) {
  219.             BGR_mean[((iter167) - (0))] = ((iter167 == 0) ? norm_mean_0 : ((iter167 == 1) ? norm_mean_1 : norm_mean_2));
  220.         }
  221.     }
  222.     const uint64_t iter187_iter188_fused_iter189_fused = ((((((blockIdx.x) * (128))) + (threadIdx.x))) + (0));
  223.     const uint64_t iter189 = (((iter187_iter188_fused_iter189_fused % 3)) + (0));
  224.     const uint64_t iter187_iter188_fused = ((((iter187_iter188_fused_iter189_fused) / (3))) + (0));
  225.     const uint64_t iter188 = (((iter187_iter188_fused % crop_w)) + (0));
  226.     const uint64_t iter187 = ((((iter187_iter188_fused) / (crop_w))) + (0));
  227.     if ((iter187 < ((0) + (crop_h)))) {
  228.         if ((iter188 < ((0) + (crop_w)))) {
  229.             if ((iter189 < ((0) + (3)))) {
  230.                 HWC2CHW[((((((iter188) - (0))) + (((crop_w) * (((iter187) - (0))))))) + (((((crop_w) * (crop_h))) * (((iter189) - (0))))))] = (((((float)(((((((((((((((cubfh[((((iter187) + (crop_top))) + (((resize_h) * (0))))]) * (cubfw[((((iter188) + (crop_left))) + (((resize_w) * (0))))]))) * (input[((((((2) - (iter189))) + (((4) * (intw[((((iter188) + (crop_left))) + (((resize_w) * (0))))]))))) + (((((4) * (w))) * (inth[((((iter187) + (crop_top))) + (((resize_h) * (0))))]))))]))) + (((((cubfh[((((iter187) + (crop_top))) + (((resize_h) * (1))))]) * (cubfw[((((iter188) + (crop_left))) + (((resize_w) * (0))))]))) * (input[((((((2) - (iter189))) + (((4) * (intw[((((iter188) + (crop_left))) + (((resize_w) * (0))))]))))) + (((((4) * (w))) * (inth[((((iter187) + (crop_top))) + (((resize_h) * (1))))]))))]))))) + (((((cubfh[((((iter187) + (crop_top))) + (((resize_h) * (0))))]) * (cubfw[((((iter188) + (crop_left))) + (((resize_w) * (1))))]))) * (input[((((((2) - (iter189))) + (((4) * (intw[((((iter188) + (crop_left))) + (((resize_w) * (1))))]))))) + (((((4) * (w))) * (inth[((((iter187) + (crop_top))) + (((resize_h) * (0))))]))))]))))) + (((((cubfh[((((iter187) + (crop_top))) + (((resize_h) * (1))))]) * (cubfw[((((iter188) + (crop_left))) + (((resize_w) * (1))))]))) * (input[((((((2) - (iter189))) + (((4) * (intw[((((iter188) + (crop_left))) + (((resize_w) * (1))))]))))) + (((((4) * (w))) * (inth[((((iter187) + (crop_top))) + (((resize_h) * (1))))]))))]))))) + (2097152))) / (4194304)))) - (BGR_mean[((iter189) - (0))]))) / (BGR_std[((iter189) - (0))]));
  231.             }
  232.         }
  233.     }
  234. }
  235.  
  236.  __global__ void BGRA_Nearest_Kernel(uint64_t resize_h, uint64_t resize_w, uint64_t crop_h, uint64_t crop_w, int32_t crop_top, int32_t crop_left, float norm_mean_0, float norm_mean_1, float norm_mean_2, float norm_std_0, float norm_std_1, float norm_std_2, uint64_t pad_h, uint64_t pad_w, int32_t pad_top, int32_t pad_left, int32_t pad_bottom, int32_t pad_right, float pad_value, uint8_t* __restrict__ input, float* __restrict__ HWC2CHW, uint64_t h, uint64_t w) {
  237.     float BGR_std[3];
  238.     float scale[2];
  239.     float BGR_mean[3];
  240.     for (uint64_t iter196 = 0; iter196 < 0 + 3; ++iter196) {
  241.         if ((iter196 < ((0) + (3)))) {
  242.             BGR_std[((iter196) - (0))] = ((iter196 == 0) ? norm_std_0 : ((iter196 == 1) ? norm_std_1 : norm_std_2));
  243.         }
  244.     }
  245.     for (uint64_t iter203 = 0; iter203 < 0 + 2; ++iter203) {
  246.         if ((iter203 < ((0) + (2)))) {
  247.             scale[((iter203) - (0))] = ((iter203 == 0) ? (((float)(h)) / ((float)(resize_h))) : (((float)(w)) / ((float)(resize_w))));
  248.         }
  249.     }
  250.     for (uint64_t iter195 = 0; iter195 < 0 + 3; ++iter195) {
  251.         if ((iter195 < ((0) + (3)))) {
  252.             BGR_mean[((iter195) - (0))] = ((iter195 == 0) ? norm_mean_0 : ((iter195 == 1) ? norm_mean_1 : norm_mean_2));
  253.         }
  254.     }
  255.     const uint64_t iter216_iter217_fused_iter218_fused = ((((((blockIdx.x) * (128))) + (threadIdx.x))) + (0));
  256.     const uint64_t iter218 = (((iter216_iter217_fused_iter218_fused % 3)) + (0));
  257.     const uint64_t iter216_iter217_fused = ((((iter216_iter217_fused_iter218_fused) / (3))) + (0));
  258.     const uint64_t iter217 = (((iter216_iter217_fused % crop_w)) + (0));
  259.     const uint64_t iter216 = ((((iter216_iter217_fused) / (crop_w))) + (0));
  260.     if ((iter216 < ((0) + (crop_h)))) {
  261.         if ((iter217 < ((0) + (crop_w)))) {
  262.             if ((iter218 < ((0) + (3)))) {
  263.                 HWC2CHW[((((((iter217) - (0))) + (((crop_w) * (((iter216) - (0))))))) + (((((crop_w) * (crop_h))) * (((iter218) - (0))))))] = (((((float)(input[((((((2) - (iter218))) + (((4) * (min((uint64_t)(static_cast<uint64_t>(floorf(((((iter217) + (crop_left))) * (scale[((1) - (0))]))))), (uint64_t)((w) - (1)))))))) + (((((4) * (w))) * (min((uint64_t)(static_cast<uint64_t>(floorf(((((iter216) + (crop_top))) * (scale[((0) - (0))]))))), (uint64_t)((h) - (1)))))))])) - (BGR_mean[((iter218) - (0))]))) / (BGR_std[((iter218) - (0))]));
  264.             }
  265.         }
  266.     }
  267. }
  268.  
  269.  __global__ void NV12_Bilinear_Kernel(uint64_t resize_h, uint64_t resize_w, uint64_t crop_h, uint64_t crop_w, int32_t crop_top, int32_t crop_left, float norm_mean_0, float norm_mean_1, float norm_mean_2, float norm_std_0, float norm_std_1, float norm_std_2, uint64_t pad_h, uint64_t pad_w, int32_t pad_top, int32_t pad_left, int32_t pad_bottom, int32_t pad_right, float pad_value, int16_t* __restrict__ cubfh, int16_t* __restrict__ cubfw, int32_t* __restrict__ inth, int32_t* __restrict__ intw, uint8_t* __restrict__ input, float* __restrict__ HWC2CHW, uint64_t h, uint64_t w) {
  270.     float BGR_std[3];
  271.     float NV2BGR_params[((3) * (3))];
  272.     float BGR_mean[3];
  273.     for (uint64_t iter225 = 0; iter225 < 0 + 3; ++iter225) {
  274.         if ((iter225 < ((0) + (3)))) {
  275.             BGR_std[((iter225) - (0))] = ((iter225 == 0) ? norm_std_0 : ((iter225 == 1) ? norm_std_1 : norm_std_2));
  276.         }
  277.     }
  278.     for (uint64_t iter222 = 0; iter222 < 0 + 3; ++iter222) {
  279.         for (uint64_t iter223 = 0; iter223 < 0 + 3; ++iter223) {
  280.             if ((iter222 < ((0) + (3)))) {
  281.                 if ((iter223 < ((0) + (3)))) {
  282.                     NV2BGR_params[((((iter223) - (0))) + (((3) * (((iter222) - (0))))))] = (((((true) & ((iter222 == 0)))) & ((iter223 == 0))) ? (static_cast<float>(1.164)) : (((((true) & ((iter222 == 0)))) & ((iter223 == 1))) ? (static_cast<float>(2.018)) : (((((true) & ((iter222 == 0)))) & ((iter223 == 2))) ? (static_cast<float>(0)) : (((((true) & ((iter222 == 1)))) & ((iter223 == 0))) ? (static_cast<float>(1.164)) : (((((true) & ((iter222 == 1)))) & ((iter223 == 1))) ? (static_cast<float>(-0.813)) : (((((true) & ((iter222 == 1)))) & ((iter223 == 2))) ? (static_cast<float>(-0.391)) : (((((true) & ((iter222 == 2)))) & ((iter223 == 0))) ? (static_cast<float>(1.164)) : (((((true) & ((iter222 == 2)))) & ((iter223 == 1))) ? (static_cast<float>(0)) : (static_cast<float>(1.596))))))))));
  283.                 }
  284.             }
  285.         }
  286.     }
  287.     for (uint64_t iter224 = 0; iter224 < 0 + 3; ++iter224) {
  288.         if ((iter224 < ((0) + (3)))) {
  289.             BGR_mean[((iter224) - (0))] = ((iter224 == 0) ? norm_mean_0 : ((iter224 == 1) ? norm_mean_1 : norm_mean_2));
  290.         }
  291.     }
  292.     const uint64_t iter244_iter245_fused_iter246_fused = ((((((blockIdx.x) * (128))) + (threadIdx.x))) + (0));
  293.     const uint64_t iter246 = (((iter244_iter245_fused_iter246_fused % 3)) + (0));
  294.     const uint64_t iter244_iter245_fused = ((((iter244_iter245_fused_iter246_fused) / (3))) + (0));
  295.     const uint64_t iter245 = (((iter244_iter245_fused % crop_w)) + (0));
  296.     const uint64_t iter244 = ((((iter244_iter245_fused) / (crop_w))) + (0));
  297.     if ((iter244 < ((0) + (crop_h)))) {
  298.         if ((iter245 < ((0) + (crop_w)))) {
  299.             if ((iter246 < ((0) + (3)))) {
  300.                 HWC2CHW[((((((iter245) - (0))) + (((crop_w) * (((iter244) - (0))))))) + (((((crop_w) * (crop_h))) * (((iter246) - (0))))))] = (((((float)(((((((((((((((cubfh[((((iter244) + (crop_top))) + (((resize_h) * (0))))]) * (cubfw[((((iter245) + (crop_left))) + (((resize_w) * (0))))]))) * (((((((NV2BGR_params[((((0) - (0))) + (((3) * (((((2) - (iter246))) - (0))))))]) * (((input[((((0) + (((1) * (intw[((((iter245) + (crop_left))) + (((resize_w) * (0))))]))))) + (((((1) * (w))) * (inth[((((iter244) + (crop_top))) + (((resize_h) * (0))))]))))]) - ((static_cast<float>(16))))))) + (((NV2BGR_params[((((1) - (0))) + (((3) * (((((2) - (iter246))) - (0))))))]) * (((input[((((0) + (((1) * (((((intw[((((iter245) + (crop_left))) + (((resize_w) * (0))))]) / (2))) * (2))))))) + (((((1) * (w))) * (((((inth[((((iter244) + (crop_top))) + (((resize_h) * (0))))]) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))) + (((NV2BGR_params[((((2) - (0))) + (((3) * (((((2) - (iter246))) - (0))))))]) * (((input[((((0) + (((1) * (((((((intw[((((iter245) + (crop_left))) + (((resize_w) * (0))))]) / (2))) * (2))) + (1))))))) + (((((1) * (w))) * (((((inth[((((iter244) + (crop_top))) + (((resize_h) * (0))))]) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))))) + (((((cubfh[((((iter244) + (crop_top))) + (((resize_h) * (1))))]) * (cubfw[((((iter245) + (crop_left))) + (((resize_w) * (0))))]))) * (((((((NV2BGR_params[((((0) - (0))) + (((3) * (((((2) - (iter246))) - (0))))))]) * (((input[((((0) + (((1) * (intw[((((iter245) + (crop_left))) + (((resize_w) * (0))))]))))) + (((((1) * (w))) * (inth[((((iter244) + (crop_top))) + (((resize_h) * (1))))]))))]) - ((static_cast<float>(16))))))) + (((NV2BGR_params[((((1) - (0))) + (((3) * (((((2) - (iter246))) - (0))))))]) * (((input[((((0) + (((1) * (((((intw[((((iter245) + (crop_left))) + (((resize_w) * (0))))]) / (2))) * (2))))))) + (((((1) * (w))) * (((((inth[((((iter244) + (crop_top))) + (((resize_h) * (1))))]) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))) + (((NV2BGR_params[((((2) - (0))) + (((3) * (((((2) - (iter246))) - (0))))))]) * (((input[((((0) + (((1) * (((((((intw[((((iter245) + (crop_left))) + (((resize_w) * (0))))]) / (2))) * (2))) + (1))))))) + (((((1) * (w))) * (((((inth[((((iter244) + (crop_top))) + (((resize_h) * (1))))]) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))))))) + (((((cubfh[((((iter244) + (crop_top))) + (((resize_h) * (0))))]) * (cubfw[((((iter245) + (crop_left))) + (((resize_w) * (1))))]))) * (((((((NV2BGR_params[((((0) - (0))) + (((3) * (((((2) - (iter246))) - (0))))))]) * (((input[((((0) + (((1) * (intw[((((iter245) + (crop_left))) + (((resize_w) * (1))))]))))) + (((((1) * (w))) * (inth[((((iter244) + (crop_top))) + (((resize_h) * (0))))]))))]) - ((static_cast<float>(16))))))) + (((NV2BGR_params[((((1) - (0))) + (((3) * (((((2) - (iter246))) - (0))))))]) * (((input[((((0) + (((1) * (((((intw[((((iter245) + (crop_left))) + (((resize_w) * (1))))]) / (2))) * (2))))))) + (((((1) * (w))) * (((((inth[((((iter244) + (crop_top))) + (((resize_h) * (0))))]) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))) + (((NV2BGR_params[((((2) - (0))) + (((3) * (((((2) - (iter246))) - (0))))))]) * (((input[((((0) + (((1) * (((((((intw[((((iter245) + (crop_left))) + (((resize_w) * (1))))]) / (2))) * (2))) + (1))))))) + (((((1) * (w))) * (((((inth[((((iter244) + (crop_top))) + (((resize_h) * (0))))]) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))))))) + (((((cubfh[((((iter244) + (crop_top))) + (((resize_h) * (1))))]) * (cubfw[((((iter245) + (crop_left))) + (((resize_w) * (1))))]))) * (((((((NV2BGR_params[((((0) - (0))) + (((3) * (((((2) - (iter246))) - (0))))))]) * (((input[((((0) + (((1) * (intw[((((iter245) + (crop_left))) + (((resize_w) * (1))))]))))) + (((((1) * (w))) * (inth[((((iter244) + (crop_top))) + (((resize_h) * (1))))]))))]) - ((static_cast<float>(16))))))) + (((NV2BGR_params[((((1) - (0))) + (((3) * (((((2) - (iter246))) - (0))))))]) * (((input[((((0) + (((1) * (((((intw[((((iter245) + (crop_left))) + (((resize_w) * (1))))]) / (2))) * (2))))))) + (((((1) * (w))) * (((((inth[((((iter244) + (crop_top))) + (((resize_h) * (1))))]) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))) + (((NV2BGR_params[((((2) - (0))) + (((3) * (((((2) - (iter246))) - (0))))))]) * (((input[((((0) + (((1) * (((((((intw[((((iter245) + (crop_left))) + (((resize_w) * (1))))]) / (2))) * (2))) + (1))))))) + (((((1) * (w))) * (((((inth[((((iter244) + (crop_top))) + (((resize_h) * (1))))]) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))))))) + (2097152))) / (4194304)))) - (BGR_mean[((iter246) - (0))]))) / (BGR_std[((iter246) - (0))]));
  301.             }
  302.         }
  303.     }
  304. }
  305.  
  306.  __global__ void NV12_Nearest_Kernel(uint64_t resize_h, uint64_t resize_w, uint64_t crop_h, uint64_t crop_w, int32_t crop_top, int32_t crop_left, float norm_mean_0, float norm_mean_1, float norm_mean_2, float norm_std_0, float norm_std_1, float norm_std_2, uint64_t pad_h, uint64_t pad_w, int32_t pad_top, int32_t pad_left, int32_t pad_bottom, int32_t pad_right, float pad_value, uint8_t* __restrict__ input, float* __restrict__ HWC2CHW, uint64_t h, uint64_t w) {
  307.     float BGR_std[3];
  308.     float scale[2];
  309.     float NV2BGR_params[((3) * (3))];
  310.     float BGR_mean[3];
  311.     for (uint64_t iter253 = 0; iter253 < 0 + 3; ++iter253) {
  312.         if ((iter253 < ((0) + (3)))) {
  313.             BGR_std[((iter253) - (0))] = ((iter253 == 0) ? norm_std_0 : ((iter253 == 1) ? norm_std_1 : norm_std_2));
  314.         }
  315.     }
  316.     for (uint64_t iter260 = 0; iter260 < 0 + 2; ++iter260) {
  317.         if ((iter260 < ((0) + (2)))) {
  318.             scale[((iter260) - (0))] = ((iter260 == 0) ? (((float)(h)) / ((float)(resize_h))) : (((float)(w)) / ((float)(resize_w))));
  319.         }
  320.     }
  321.     for (uint64_t iter250 = 0; iter250 < 0 + 3; ++iter250) {
  322.         for (uint64_t iter251 = 0; iter251 < 0 + 3; ++iter251) {
  323.             if ((iter250 < ((0) + (3)))) {
  324.                 if ((iter251 < ((0) + (3)))) {
  325.                     NV2BGR_params[((((iter251) - (0))) + (((3) * (((iter250) - (0))))))] = (((((true) & ((iter250 == 0)))) & ((iter251 == 0))) ? (static_cast<float>(1.164)) : (((((true) & ((iter250 == 0)))) & ((iter251 == 1))) ? (static_cast<float>(2.018)) : (((((true) & ((iter250 == 0)))) & ((iter251 == 2))) ? (static_cast<float>(0)) : (((((true) & ((iter250 == 1)))) & ((iter251 == 0))) ? (static_cast<float>(1.164)) : (((((true) & ((iter250 == 1)))) & ((iter251 == 1))) ? (static_cast<float>(-0.813)) : (((((true) & ((iter250 == 1)))) & ((iter251 == 2))) ? (static_cast<float>(-0.391)) : (((((true) & ((iter250 == 2)))) & ((iter251 == 0))) ? (static_cast<float>(1.164)) : (((((true) & ((iter250 == 2)))) & ((iter251 == 1))) ? (static_cast<float>(0)) : (static_cast<float>(1.596))))))))));
  326.                 }
  327.             }
  328.         }
  329.     }
  330.     for (uint64_t iter252 = 0; iter252 < 0 + 3; ++iter252) {
  331.         if ((iter252 < ((0) + (3)))) {
  332.             BGR_mean[((iter252) - (0))] = ((iter252 == 0) ? norm_mean_0 : ((iter252 == 1) ? norm_mean_1 : norm_mean_2));
  333.         }
  334.     }
  335.     const uint64_t iter273_iter274_fused_iter275_fused = ((((((blockIdx.x) * (128))) + (threadIdx.x))) + (0));
  336.     const uint64_t iter275 = (((iter273_iter274_fused_iter275_fused % 3)) + (0));
  337.     const uint64_t iter273_iter274_fused = ((((iter273_iter274_fused_iter275_fused) / (3))) + (0));
  338.     const uint64_t iter274 = (((iter273_iter274_fused % crop_w)) + (0));
  339.     const uint64_t iter273 = ((((iter273_iter274_fused) / (crop_w))) + (0));
  340.     if ((iter273 < ((0) + (crop_h)))) {
  341.         if ((iter274 < ((0) + (crop_w)))) {
  342.             if ((iter275 < ((0) + (3)))) {
  343.                 HWC2CHW[((((((iter274) - (0))) + (((crop_w) * (((iter273) - (0))))))) + (((((crop_w) * (crop_h))) * (((iter275) - (0))))))] = (((((float)(((((((NV2BGR_params[((((0) - (0))) + (((3) * (((((2) - (iter275))) - (0))))))]) * (((input[((((0) + (((1) * (min((uint64_t)(static_cast<uint64_t>(floorf(((((iter274) + (crop_left))) * (scale[((1) - (0))]))))), (uint64_t)((w) - (1)))))))) + (((((1) * (w))) * (min((uint64_t)(static_cast<uint64_t>(floorf(((((iter273) + (crop_top))) * (scale[((0) - (0))]))))), (uint64_t)((h) - (1)))))))]) - ((static_cast<float>(16))))))) + (((NV2BGR_params[((((1) - (0))) + (((3) * (((((2) - (iter275))) - (0))))))]) * (((input[((((0) + (((1) * (((((min((uint64_t)(static_cast<uint64_t>(floorf(((((iter274) + (crop_left))) * (scale[((1) - (0))]))))), (uint64_t)((w) - (1)))) / (2))) * (2))))))) + (((((1) * (w))) * (((((min((uint64_t)(static_cast<uint64_t>(floorf(((((iter273) + (crop_top))) * (scale[((0) - (0))]))))), (uint64_t)((h) - (1)))) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))) + (((NV2BGR_params[((((2) - (0))) + (((3) * (((((2) - (iter275))) - (0))))))]) * (((input[((((0) + (((1) * (((((((min((uint64_t)(static_cast<uint64_t>(floorf(((((iter274) + (crop_left))) * (scale[((1) - (0))]))))), (uint64_t)((w) - (1)))) / (2))) * (2))) + (1))))))) + (((((1) * (w))) * (((((min((uint64_t)(static_cast<uint64_t>(floorf(((((iter273) + (crop_top))) * (scale[((0) - (0))]))))), (uint64_t)((h) - (1)))) / (2))) + (h))))))]) - ((static_cast<float>(128)))))))))) - (BGR_mean[((iter275) - (0))]))) / (BGR_std[((iter275) - (0))]));
  344.             }
  345.         }
  346.     }
  347. }
  348.  
  349.  __global__ void NV21_Bilinear_Kernel(uint64_t resize_h, uint64_t resize_w, uint64_t crop_h, uint64_t crop_w, int32_t crop_top, int32_t crop_left, float norm_mean_0, float norm_mean_1, float norm_mean_2, float norm_std_0, float norm_std_1, float norm_std_2, uint64_t pad_h, uint64_t pad_w, int32_t pad_top, int32_t pad_left, int32_t pad_bottom, int32_t pad_right, float pad_value, int16_t* __restrict__ cubfh, int16_t* __restrict__ cubfw, int32_t* __restrict__ inth, int32_t* __restrict__ intw, uint8_t* __restrict__ input, float* __restrict__ HWC2CHW, uint64_t h, uint64_t w) {
  350.     float BGR_std[3];
  351.     float NV2BGR_params[((3) * (3))];
  352.     float BGR_mean[3];
  353.     for (uint64_t iter282 = 0; iter282 < 0 + 3; ++iter282) {
  354.         if ((iter282 < ((0) + (3)))) {
  355.             BGR_std[((iter282) - (0))] = ((iter282 == 0) ? norm_std_0 : ((iter282 == 1) ? norm_std_1 : norm_std_2));
  356.         }
  357.     }
  358.     for (uint64_t iter279 = 0; iter279 < 0 + 3; ++iter279) {
  359.         for (uint64_t iter280 = 0; iter280 < 0 + 3; ++iter280) {
  360.             if ((iter279 < ((0) + (3)))) {
  361.                 if ((iter280 < ((0) + (3)))) {
  362.                     NV2BGR_params[((((iter280) - (0))) + (((3) * (((iter279) - (0))))))] = (((((true) & ((iter279 == 0)))) & ((iter280 == 0))) ? (static_cast<float>(1.164)) : (((((true) & ((iter279 == 0)))) & ((iter280 == 1))) ? (static_cast<float>(2.018)) : (((((true) & ((iter279 == 0)))) & ((iter280 == 2))) ? (static_cast<float>(0)) : (((((true) & ((iter279 == 1)))) & ((iter280 == 0))) ? (static_cast<float>(1.164)) : (((((true) & ((iter279 == 1)))) & ((iter280 == 1))) ? (static_cast<float>(-0.813)) : (((((true) & ((iter279 == 1)))) & ((iter280 == 2))) ? (static_cast<float>(-0.391)) : (((((true) & ((iter279 == 2)))) & ((iter280 == 0))) ? (static_cast<float>(1.164)) : (((((true) & ((iter279 == 2)))) & ((iter280 == 1))) ? (static_cast<float>(0)) : (static_cast<float>(1.596))))))))));
  363.                 }
  364.             }
  365.         }
  366.     }
  367.     for (uint64_t iter281 = 0; iter281 < 0 + 3; ++iter281) {
  368.         if ((iter281 < ((0) + (3)))) {
  369.             BGR_mean[((iter281) - (0))] = ((iter281 == 0) ? norm_mean_0 : ((iter281 == 1) ? norm_mean_1 : norm_mean_2));
  370.         }
  371.     }
  372.     const uint64_t iter301_iter302_fused_iter303_fused = ((((((blockIdx.x) * (128))) + (threadIdx.x))) + (0));
  373.     const uint64_t iter303 = (((iter301_iter302_fused_iter303_fused % 3)) + (0));
  374.     const uint64_t iter301_iter302_fused = ((((iter301_iter302_fused_iter303_fused) / (3))) + (0));
  375.     const uint64_t iter302 = (((iter301_iter302_fused % crop_w)) + (0));
  376.     const uint64_t iter301 = ((((iter301_iter302_fused) / (crop_w))) + (0));
  377.     if ((iter301 < ((0) + (crop_h)))) {
  378.         if ((iter302 < ((0) + (crop_w)))) {
  379.             if ((iter303 < ((0) + (3)))) {
  380.                 HWC2CHW[((((((iter302) - (0))) + (((crop_w) * (((iter301) - (0))))))) + (((((crop_w) * (crop_h))) * (((iter303) - (0))))))] = (((((float)(((((((((((((((cubfh[((((iter301) + (crop_top))) + (((resize_h) * (0))))]) * (cubfw[((((iter302) + (crop_left))) + (((resize_w) * (0))))]))) * (((((((NV2BGR_params[((((0) - (0))) + (((3) * (((((2) - (iter303))) - (0))))))]) * (((input[((((0) + (((1) * (intw[((((iter302) + (crop_left))) + (((resize_w) * (0))))]))))) + (((((1) * (w))) * (inth[((((iter301) + (crop_top))) + (((resize_h) * (0))))]))))]) - ((static_cast<float>(16))))))) + (((NV2BGR_params[((((1) - (0))) + (((3) * (((((2) - (iter303))) - (0))))))]) * (((input[((((0) + (((1) * (((((((intw[((((iter302) + (crop_left))) + (((resize_w) * (0))))]) / (2))) * (2))) + (1))))))) + (((((1) * (w))) * (((((inth[((((iter301) + (crop_top))) + (((resize_h) * (0))))]) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))) + (((NV2BGR_params[((((2) - (0))) + (((3) * (((((2) - (iter303))) - (0))))))]) * (((input[((((0) + (((1) * (((((intw[((((iter302) + (crop_left))) + (((resize_w) * (0))))]) / (2))) * (2))))))) + (((((1) * (w))) * (((((inth[((((iter301) + (crop_top))) + (((resize_h) * (0))))]) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))))) + (((((cubfh[((((iter301) + (crop_top))) + (((resize_h) * (1))))]) * (cubfw[((((iter302) + (crop_left))) + (((resize_w) * (0))))]))) * (((((((NV2BGR_params[((((0) - (0))) + (((3) * (((((2) - (iter303))) - (0))))))]) * (((input[((((0) + (((1) * (intw[((((iter302) + (crop_left))) + (((resize_w) * (0))))]))))) + (((((1) * (w))) * (inth[((((iter301) + (crop_top))) + (((resize_h) * (1))))]))))]) - ((static_cast<float>(16))))))) + (((NV2BGR_params[((((1) - (0))) + (((3) * (((((2) - (iter303))) - (0))))))]) * (((input[((((0) + (((1) * (((((((intw[((((iter302) + (crop_left))) + (((resize_w) * (0))))]) / (2))) * (2))) + (1))))))) + (((((1) * (w))) * (((((inth[((((iter301) + (crop_top))) + (((resize_h) * (1))))]) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))) + (((NV2BGR_params[((((2) - (0))) + (((3) * (((((2) - (iter303))) - (0))))))]) * (((input[((((0) + (((1) * (((((intw[((((iter302) + (crop_left))) + (((resize_w) * (0))))]) / (2))) * (2))))))) + (((((1) * (w))) * (((((inth[((((iter301) + (crop_top))) + (((resize_h) * (1))))]) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))))))) + (((((cubfh[((((iter301) + (crop_top))) + (((resize_h) * (0))))]) * (cubfw[((((iter302) + (crop_left))) + (((resize_w) * (1))))]))) * (((((((NV2BGR_params[((((0) - (0))) + (((3) * (((((2) - (iter303))) - (0))))))]) * (((input[((((0) + (((1) * (intw[((((iter302) + (crop_left))) + (((resize_w) * (1))))]))))) + (((((1) * (w))) * (inth[((((iter301) + (crop_top))) + (((resize_h) * (0))))]))))]) - ((static_cast<float>(16))))))) + (((NV2BGR_params[((((1) - (0))) + (((3) * (((((2) - (iter303))) - (0))))))]) * (((input[((((0) + (((1) * (((((((intw[((((iter302) + (crop_left))) + (((resize_w) * (1))))]) / (2))) * (2))) + (1))))))) + (((((1) * (w))) * (((((inth[((((iter301) + (crop_top))) + (((resize_h) * (0))))]) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))) + (((NV2BGR_params[((((2) - (0))) + (((3) * (((((2) - (iter303))) - (0))))))]) * (((input[((((0) + (((1) * (((((intw[((((iter302) + (crop_left))) + (((resize_w) * (1))))]) / (2))) * (2))))))) + (((((1) * (w))) * (((((inth[((((iter301) + (crop_top))) + (((resize_h) * (0))))]) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))))))) + (((((cubfh[((((iter301) + (crop_top))) + (((resize_h) * (1))))]) * (cubfw[((((iter302) + (crop_left))) + (((resize_w) * (1))))]))) * (((((((NV2BGR_params[((((0) - (0))) + (((3) * (((((2) - (iter303))) - (0))))))]) * (((input[((((0) + (((1) * (intw[((((iter302) + (crop_left))) + (((resize_w) * (1))))]))))) + (((((1) * (w))) * (inth[((((iter301) + (crop_top))) + (((resize_h) * (1))))]))))]) - ((static_cast<float>(16))))))) + (((NV2BGR_params[((((1) - (0))) + (((3) * (((((2) - (iter303))) - (0))))))]) * (((input[((((0) + (((1) * (((((((intw[((((iter302) + (crop_left))) + (((resize_w) * (1))))]) / (2))) * (2))) + (1))))))) + (((((1) * (w))) * (((((inth[((((iter301) + (crop_top))) + (((resize_h) * (1))))]) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))) + (((NV2BGR_params[((((2) - (0))) + (((3) * (((((2) - (iter303))) - (0))))))]) * (((input[((((0) + (((1) * (((((intw[((((iter302) + (crop_left))) + (((resize_w) * (1))))]) / (2))) * (2))))))) + (((((1) * (w))) * (((((inth[((((iter301) + (crop_top))) + (((resize_h) * (1))))]) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))))))) + (2097152))) / (4194304)))) - (BGR_mean[((iter303) - (0))]))) / (BGR_std[((iter303) - (0))]));
  381.             }
  382.         }
  383.     }
  384. }
  385.  
  386.  __global__ void NV21_Nearest_Kernel(uint64_t resize_h, uint64_t resize_w, uint64_t crop_h, uint64_t crop_w, int32_t crop_top, int32_t crop_left, float norm_mean_0, float norm_mean_1, float norm_mean_2, float norm_std_0, float norm_std_1, float norm_std_2, uint64_t pad_h, uint64_t pad_w, int32_t pad_top, int32_t pad_left, int32_t pad_bottom, int32_t pad_right, float pad_value, uint8_t* __restrict__ input, float* __restrict__ HWC2CHW, uint64_t h, uint64_t w) {
  387.     float BGR_std[3];
  388.     float scale[2];
  389.     float NV2BGR_params[((3) * (3))];
  390.     float BGR_mean[3];
  391.     for (uint64_t iter310 = 0; iter310 < 0 + 3; ++iter310) {
  392.         if ((iter310 < ((0) + (3)))) {
  393.             BGR_std[((iter310) - (0))] = ((iter310 == 0) ? norm_std_0 : ((iter310 == 1) ? norm_std_1 : norm_std_2));
  394.         }
  395.     }
  396.     for (uint64_t iter317 = 0; iter317 < 0 + 2; ++iter317) {
  397.         if ((iter317 < ((0) + (2)))) {
  398.             scale[((iter317) - (0))] = ((iter317 == 0) ? (((float)(h)) / ((float)(resize_h))) : (((float)(w)) / ((float)(resize_w))));
  399.         }
  400.     }
  401.     for (uint64_t iter307 = 0; iter307 < 0 + 3; ++iter307) {
  402.         for (uint64_t iter308 = 0; iter308 < 0 + 3; ++iter308) {
  403.             if ((iter307 < ((0) + (3)))) {
  404.                 if ((iter308 < ((0) + (3)))) {
  405.                     NV2BGR_params[((((iter308) - (0))) + (((3) * (((iter307) - (0))))))] = (((((true) & ((iter307 == 0)))) & ((iter308 == 0))) ? (static_cast<float>(1.164)) : (((((true) & ((iter307 == 0)))) & ((iter308 == 1))) ? (static_cast<float>(2.018)) : (((((true) & ((iter307 == 0)))) & ((iter308 == 2))) ? (static_cast<float>(0)) : (((((true) & ((iter307 == 1)))) & ((iter308 == 0))) ? (static_cast<float>(1.164)) : (((((true) & ((iter307 == 1)))) & ((iter308 == 1))) ? (static_cast<float>(-0.813)) : (((((true) & ((iter307 == 1)))) & ((iter308 == 2))) ? (static_cast<float>(-0.391)) : (((((true) & ((iter307 == 2)))) & ((iter308 == 0))) ? (static_cast<float>(1.164)) : (((((true) & ((iter307 == 2)))) & ((iter308 == 1))) ? (static_cast<float>(0)) : (static_cast<float>(1.596))))))))));
  406.                 }
  407.             }
  408.         }
  409.     }
  410.     for (uint64_t iter309 = 0; iter309 < 0 + 3; ++iter309) {
  411.         if ((iter309 < ((0) + (3)))) {
  412.             BGR_mean[((iter309) - (0))] = ((iter309 == 0) ? norm_mean_0 : ((iter309 == 1) ? norm_mean_1 : norm_mean_2));
  413.         }
  414.     }
  415.     const uint64_t iter330_iter331_fused_iter332_fused = ((((((blockIdx.x) * (128))) + (threadIdx.x))) + (0));
  416.     const uint64_t iter332 = (((iter330_iter331_fused_iter332_fused % 3)) + (0));
  417.     const uint64_t iter330_iter331_fused = ((((iter330_iter331_fused_iter332_fused) / (3))) + (0));
  418.     const uint64_t iter331 = (((iter330_iter331_fused % crop_w)) + (0));
  419.     const uint64_t iter330 = ((((iter330_iter331_fused) / (crop_w))) + (0));
  420.     if ((iter330 < ((0) + (crop_h)))) {
  421.         if ((iter331 < ((0) + (crop_w)))) {
  422.             if ((iter332 < ((0) + (3)))) {
  423.                 HWC2CHW[((((((iter331) - (0))) + (((crop_w) * (((iter330) - (0))))))) + (((((crop_w) * (crop_h))) * (((iter332) - (0))))))] = (((((float)(((((((NV2BGR_params[((((0) - (0))) + (((3) * (((((2) - (iter332))) - (0))))))]) * (((input[((((0) + (((1) * (min((uint64_t)(static_cast<uint64_t>(floorf(((((iter331) + (crop_left))) * (scale[((1) - (0))]))))), (uint64_t)((w) - (1)))))))) + (((((1) * (w))) * (min((uint64_t)(static_cast<uint64_t>(floorf(((((iter330) + (crop_top))) * (scale[((0) - (0))]))))), (uint64_t)((h) - (1)))))))]) - ((static_cast<float>(16))))))) + (((NV2BGR_params[((((1) - (0))) + (((3) * (((((2) - (iter332))) - (0))))))]) * (((input[((((0) + (((1) * (((((((min((uint64_t)(static_cast<uint64_t>(floorf(((((iter331) + (crop_left))) * (scale[((1) - (0))]))))), (uint64_t)((w) - (1)))) / (2))) * (2))) + (1))))))) + (((((1) * (w))) * (((((min((uint64_t)(static_cast<uint64_t>(floorf(((((iter330) + (crop_top))) * (scale[((0) - (0))]))))), (uint64_t)((h) - (1)))) / (2))) + (h))))))]) - ((static_cast<float>(128))))))))) + (((NV2BGR_params[((((2) - (0))) + (((3) * (((((2) - (iter332))) - (0))))))]) * (((input[((((0) + (((1) * (((((min((uint64_t)(static_cast<uint64_t>(floorf(((((iter331) + (crop_left))) * (scale[((1) - (0))]))))), (uint64_t)((w) - (1)))) / (2))) * (2))))))) + (((((1) * (w))) * (((((min((uint64_t)(static_cast<uint64_t>(floorf(((((iter330) + (crop_top))) * (scale[((0) - (0))]))))), (uint64_t)((h) - (1)))) / (2))) + (h))))))]) - ((static_cast<float>(128)))))))))) - (BGR_mean[((iter332) - (0))]))) / (BGR_std[((iter332) - (0))]));
  424.             }
  425.         }
  426.     }
  427. }
  428.  
  429.  
  430. #define INCREASE(x, l) ((x + 1) >= (l) ? (x) : ((x) + 1))
  431.  
  432.  __global__ void bilinear_resize_preprocess_h(uint64_t src_h, uint64_t dst_h,
  433.     int16_t* __restrict__ cubfh, int32_t* __restrict__ inth) {
  434.  
  435.     int element_x = blockIdx.x * blockDim.x + threadIdx.x;
  436.     if (element_x >= dst_h) {
  437.         return;
  438.     }
  439.  
  440.     float scale_h = double(src_h) / dst_h;
  441.  
  442.     float fh = (float)((element_x + 0.5) * scale_h - 0.5f);
  443.     int sh = floor(fh);
  444.     fh -= sh;
  445.     if (sh < 0) {
  446.       fh = 0;
  447.       sh = 0;
  448.     }
  449.     if (sh >= src_h) {
  450.       fh = 0;
  451.       sh = src_h - 1;
  452.     }
  453.  
  454.     int int_h1 = INCREASE(sh, src_h);
  455.  
  456.     fh = fh * 2048;
  457.     cubfh[element_x] = rint(2048 - fh);
  458.     cubfh[dst_h + element_x] = rint(fh);
  459.  
  460.     inth[element_x] = sh;
  461.     inth[dst_h + element_x] = int_h1;
  462. }
  463.  
  464.  __global__ void bilinear_resize_preprocess_w(uint64_t src_w, uint64_t dst_w,
  465.     int16_t* __restrict__ cubfw, int32_t* __restrict__ intw) {
  466.  
  467.     int element_x = blockIdx.x * blockDim.x + threadIdx.x;
  468.     if (element_x >= dst_w) {
  469.         return;
  470.     }
  471.  
  472.     float scale_w = double(src_w) / dst_w;
  473.  
  474.     float fw = (float)((element_x + 0.5) * scale_w - 0.5f);
  475.     int sw = floor(fw);
  476.     fw -= sw;
  477.     if (sw < 0) {
  478.       fw = 0;
  479.       sw = 0;
  480.     }
  481.     if (sw >= src_w) {
  482.       fw = 0;
  483.       sw = src_w - 1;
  484.     }
  485.  
  486.     int int_w1 = INCREASE(sw, src_w);
  487.  
  488.     fw = fw * 2048;
  489.     cubfw[element_x] = rint(2048 - fw);
  490.     cubfw[dst_w + element_x] = rint(fw);
  491.  
  492.     intw[element_x] = sw;
  493.     intw[dst_w + element_x] = int_w1;
  494. }
  495.  
  496.  
  497.  
  498.  void FuseKernelCU(cudaStream_t stream, uint64_t resize_h, uint64_t resize_w, uint64_t crop_h, uint64_t crop_w, int32_t crop_top, int32_t crop_left, float norm_mean_0, float norm_mean_1, float norm_mean_2, float norm_std_0, float norm_std_1, float norm_std_2, uint64_t pad_h, uint64_t pad_w, int32_t pad_top, int32_t pad_left, int32_t pad_bottom, int32_t pad_right, float pad_value, uint8_t* __restrict__ src_raw_data, float* __restrict__ dst_raw_data, uint64_t dst_element_num, uint64_t src_h, uint64_t src_w, const char *format, const char *interpolation = "nearest"){
  499.  
  500.     if (resize_h && resize_w && EQUAL(interpolation, "nearest")) {
  501.         if(EQUAL(format, "BGR")){
  502.           BGR_Nearest_Kernel<<<(dst_element_num + BLOCK_SIZE -1) / BLOCK_SIZE, BLOCK_SIZE, 0, stream>>>(resize_h, resize_w, crop_h, crop_w, crop_top, crop_left, norm_mean_0, norm_mean_1, norm_mean_2, norm_std_0, norm_std_1, norm_std_2, pad_h, pad_w, pad_top, pad_left, pad_bottom, pad_right, pad_value, src_raw_data, dst_raw_data, src_h, src_w);
  503.         } else if(EQUAL(format, "RGB")){
  504.           RGB_Nearest_Kernel<<<(dst_element_num + BLOCK_SIZE -1) / BLOCK_SIZE, BLOCK_SIZE, 0, stream>>>(resize_h, resize_w, crop_h, crop_w, crop_top, crop_left, norm_mean_0, norm_mean_1, norm_mean_2, norm_std_0, norm_std_1, norm_std_2, pad_h, pad_w, pad_top, pad_left, pad_bottom, pad_right, pad_value, src_raw_data, dst_raw_data, src_h, src_w);
  505.         } else if(EQUAL(format, "GRAY")){
  506.           GRAY_Nearest_Kernel<<<(dst_element_num + BLOCK_SIZE -1) / BLOCK_SIZE, BLOCK_SIZE, 0, stream>>>(resize_h, resize_w, crop_h, crop_w, crop_top, crop_left, norm_mean_0, norm_mean_1, norm_mean_2, norm_std_0, norm_std_1, norm_std_2, pad_h, pad_w, pad_top, pad_left, pad_bottom, pad_right, pad_value, src_raw_data, dst_raw_data, src_h, src_w);
  507.         } else if(EQUAL(format, "BGRA")){
  508.           BGRA_Nearest_Kernel<<<(dst_element_num + BLOCK_SIZE -1) / BLOCK_SIZE, BLOCK_SIZE, 0, stream>>>(resize_h, resize_w, crop_h, crop_w, crop_top, crop_left, norm_mean_0, norm_mean_1, norm_mean_2, norm_std_0, norm_std_1, norm_std_2, pad_h, pad_w, pad_top, pad_left, pad_bottom, pad_right, pad_value, src_raw_data, dst_raw_data, src_h, src_w);
  509.         } else if(EQUAL(format, "NV12")){
  510.           NV12_Nearest_Kernel<<<(dst_element_num + BLOCK_SIZE -1) / BLOCK_SIZE, BLOCK_SIZE, 0, stream>>>(resize_h, resize_w, crop_h, crop_w, crop_top, crop_left, norm_mean_0, norm_mean_1, norm_mean_2, norm_std_0, norm_std_1, norm_std_2, pad_h, pad_w, pad_top, pad_left, pad_bottom, pad_right, pad_value, src_raw_data, dst_raw_data, src_h, src_w);
  511.         } else if(EQUAL(format, "NV21")){
  512.           NV21_Nearest_Kernel<<<(dst_element_num + BLOCK_SIZE -1) / BLOCK_SIZE, BLOCK_SIZE, 0, stream>>>(resize_h, resize_w, crop_h, crop_w, crop_top, crop_left, norm_mean_0, norm_mean_1, norm_mean_2, norm_std_0, norm_std_1, norm_std_2, pad_h, pad_w, pad_top, pad_left, pad_bottom, pad_right, pad_value, src_raw_data, dst_raw_data, src_h, src_w);
  513.         } else {
  514.            ABORT("This format is not supported");
  515.         }
  516.     } 
  517.     else if(resize_h && resize_w && EQUAL(interpolation, "bilinear")){
  518.         short* cubfh;
  519.         short* cubfw;
  520.         int* inth;
  521.         int* intw;
  522.  
  523.         cuErrCheck(cudaMalloc(&cubfh, resize_h*2 * sizeof(short)));
  524.         cuErrCheck(cudaMalloc(&cubfw, resize_w*2 * sizeof(short)));
  525.         cuErrCheck(cudaMalloc(&inth, resize_h*2 * sizeof(int)));
  526.         cuErrCheck(cudaMalloc(&intw, resize_w*2 * sizeof(int)));
  527.  
  528.         int block = 512;
  529.         bilinear_resize_preprocess_h<<<(resize_h + block -1) / block, block, 0, stream>>>(src_h, resize_h, cubfh, inth);
  530.         bilinear_resize_preprocess_w<<<(resize_w + block -1) / block, block, 0, stream>>>(src_w, resize_w, cubfw, intw);
  531.  
  532.         if(EQUAL(format, "BGR")){
  533.           BGR_Bilinear_Kernel<<<(dst_element_num + BLOCK_SIZE -1) / BLOCK_SIZE, BLOCK_SIZE, 0, stream>>>(resize_h, resize_w, crop_h, crop_w, crop_top, crop_left, norm_mean_0, norm_mean_1, norm_mean_2, norm_std_0, norm_std_1, norm_std_2, pad_h, pad_w, pad_top, pad_left, pad_bottom, pad_right, pad_value, cubfh, cubfw, inth, intw, src_raw_data, dst_raw_data, src_h, src_w);
  534.         } else if(EQUAL(format, "RGB")){
  535.           RGB_Bilinear_Kernel<<<(dst_element_num + BLOCK_SIZE -1) / BLOCK_SIZE, BLOCK_SIZE, 0, stream>>>(resize_h, resize_w, crop_h, crop_w, crop_top, crop_left, norm_mean_0, norm_mean_1, norm_mean_2, norm_std_0, norm_std_1, norm_std_2, pad_h, pad_w, pad_top, pad_left, pad_bottom, pad_right, pad_value, cubfh, cubfw, inth, intw, src_raw_data, dst_raw_data, src_h, src_w);
  536.         } else if(EQUAL(format, "GRAY")){
  537.           GRAY_Bilinear_Kernel<<<(dst_element_num + BLOCK_SIZE -1) / BLOCK_SIZE, BLOCK_SIZE, 0, stream>>>(resize_h, resize_w, crop_h, crop_w, crop_top, crop_left, norm_mean_0, norm_mean_1, norm_mean_2, norm_std_0, norm_std_1, norm_std_2, pad_h, pad_w, pad_top, pad_left, pad_bottom, pad_right, pad_value, cubfh, cubfw, inth, intw, src_raw_data, dst_raw_data, src_h, src_w);
  538.         } else if(EQUAL(format, "BGRA")){
  539.           BGRA_Bilinear_Kernel<<<(dst_element_num + BLOCK_SIZE -1) / BLOCK_SIZE, BLOCK_SIZE, 0, stream>>>(resize_h, resize_w, crop_h, crop_w, crop_top, crop_left, norm_mean_0, norm_mean_1, norm_mean_2, norm_std_0, norm_std_1, norm_std_2, pad_h, pad_w, pad_top, pad_left, pad_bottom, pad_right, pad_value, cubfh, cubfw, inth, intw, src_raw_data, dst_raw_data, src_h, src_w);
  540.         } else if(EQUAL(format, "NV12")){
  541.           NV12_Bilinear_Kernel<<<(dst_element_num + BLOCK_SIZE -1) / BLOCK_SIZE, BLOCK_SIZE, 0, stream>>>(resize_h, resize_w, crop_h, crop_w, crop_top, crop_left, norm_mean_0, norm_mean_1, norm_mean_2, norm_std_0, norm_std_1, norm_std_2, pad_h, pad_w, pad_top, pad_left, pad_bottom, pad_right, pad_value, cubfh, cubfw, inth, intw, src_raw_data, dst_raw_data, src_h, src_w);
  542.         } else if(EQUAL(format, "NV21")){
  543.           NV21_Bilinear_Kernel<<<(dst_element_num + BLOCK_SIZE -1) / BLOCK_SIZE, BLOCK_SIZE, 0, stream>>>(resize_h, resize_w, crop_h, crop_w, crop_top, crop_left, norm_mean_0, norm_mean_1, norm_mean_2, norm_std_0, norm_std_1, norm_std_2, pad_h, pad_w, pad_top, pad_left, pad_bottom, pad_right, pad_value, cubfh, cubfw, inth, intw, src_raw_data, dst_raw_data, src_h, src_w);
  544.         } else {
  545.            ABORT("This format is not supported");
  546.         }
  547.  
  548.         cudaStreamSynchronize(stream);
  549.  
  550.         if (cubfh) cuErrCheck(cudaFree(cubfh));
  551.         if (cubfw) cuErrCheck(cudaFree(cubfw));
  552.         if (inth) cuErrCheck(cudaFree(inth));
  553.         if (intw) cuErrCheck(cudaFree(intw));
  554.     }
  555.     else {
  556.       ABORT("This interpolation is not supported");
  557.     }
  558. }
  559.  
  560.  
  561. void FuseFunc(void* stream, uint8_t* data_in, int src_h, int src_w, const char* format,
  562.               int resize_h, int resize_w, const char* interpolation, int crop_top, int crop_left,
  563.               int crop_h, int crop_w, float mean0, float mean1, float mean2, float std0, float std1,
  564.               float std2, int pad_top, int pad_left, int pad_bottom, int pad_right, int pad_h,
  565.               int pad_w, float pad_value, float* data_out, int data_out_num) {
  566.   cudaStream_t stream_ = (cudaStream_t)stream;
  567.   const char* interpolation_ = "nearest";
  568.   if (strcmp(interpolation, "bilinear") == 0) {
  569.     interpolation_ = "bilinear";
  570.   }
  571.  
  572.   FuseKernelCU(stream_, resize_h, resize_w, crop_h, crop_w, crop_top, crop_left, mean0, mean1, mean2, std0,
  573.                std1, std2, pad_top, pad_left, pad_bottom, pad_right, pad_h, pad_w, pad_value, data_in,
  574.                data_out, data_out_num, src_h, src_w, format, interpolation_);
  575. }
  576.  
  577. REGISTER_FUSE_KERNEL(fa993066c1ec32a14d8d06ecfb8d5c54e7a04ec31ba833e1abf1f16c05e027c7_cuda, "fa993066c1ec32a14d8d06ecfb8d5c54e7a04ec31ba833e1abf1f16c05e027c7_cuda",
  578.                      FuseFunc);
  579. }

Editor

You can edit this paste and save as new:


File Description
  • elena cuda
  • Paste Code
  • 12 Jul-2022
  • 56 Kb
You can Share it: