OpenI
/
MegEngine
Not watched
1
0
Fork 0
Code
Releases 31 Wiki Activity
Issues 0 Pull Requests 0 Datasets Model Cloudbrain
You can not select more than 25 topics Topics must start with a chinese character,a letter or number, can include dashes ('-') and can be up to 35 characters long.
Tree: 0a0e4b60dd
Branches Tags
${ item.name }
Create branch ${ searchTerm }
from '0a0e4b60dd'
${ noResults }
MegEngine/dnn/test/cuda/conv_bias_int8.cpp

1244 lines
55 kB
Raw Blame History 

  1. /**

  2.  * \file dnn/test/cuda/conv_bias_int8.cpp

  3.  * MegEngine is Licensed under the Apache License, Version 2.0 (the "License")

  4.  *

  5.  * Copyright (c) 2014-2020 Megvii Inc. All rights reserved.

  6.  *

  7.  * Unless required by applicable law or agreed to in writing,

  8.  * software distributed under the License is distributed on an

  9.  * "AS IS" BASIS, WITHOUT ARRANTIES OR CONDITIONS OF ANY KIND, either express or

  10.  * implied.

  11.  */

  12. #include "megdnn/oprs/nn.h"

  13. 

  14. #include "src/common/utils.h"

  15. #include "src/cuda/cudnn_with_check.h"

  16. #include "test/common/checker.h"

  17. #include "test/common/conv_bias.h"

  18. #include "test/cuda/benchmark.h"

  19. #include "test/cuda/fixture.h"

  20. #include "test/cuda/utils.h"

  21. 

  22. #define V1(x) #x

  23. #define V(x) V1(x)

  24. 

  25. namespace megdnn {

  26. namespace test {

  27. namespace {

  28. 

  29. #if MEGDNN_WITH_BENCHMARK

  30. struct BenchArgs {

  31.     size_t n, ci, hi, wi, co, f, s;

  32. };

  33. 

  34. std::vector<BenchArgs> get_resnet50_bench_args(size_t batch = 64) {

  35.     std::vector<BenchArgs> args;

  36.     args.emplace_back(BenchArgs{batch, 64, 56, 56, 256, 1, 1});

  37. 

  38.     args.emplace_back(BenchArgs{batch, 256, 56, 56, 32, 3, 1});

  39.     args.emplace_back(BenchArgs{batch, 256, 56, 56, 32, 3, 2});

  40.     args.emplace_back(BenchArgs{batch, 4, 256, 256, 32, 7, 2});

  41. 

  42.     args.emplace_back(BenchArgs{batch, 256, 56, 56, 64, 1, 1});

  43.     args.emplace_back(BenchArgs{batch, 64, 56, 56, 64, 1, 1});

  44.     args.emplace_back(BenchArgs{batch, 64, 56, 56, 64, 3, 1});

  45.     args.emplace_back(BenchArgs{batch, 64, 56, 56, 64, 3, 2});

  46.     args.emplace_back(BenchArgs{batch, 256, 56, 56, 64, 3, 2});

  47.     args.emplace_back(BenchArgs{batch, 64, 56, 56, 256, 1, 1});

  48. 

  49.     args.emplace_back(BenchArgs{batch, 256, 56, 56, 512, 1, 2});

  50.     args.emplace_back(BenchArgs{batch, 256, 56, 56, 128, 1, 2});

  51.     args.emplace_back(BenchArgs{batch, 512, 28, 28, 128, 1, 1});

  52.     args.emplace_back(BenchArgs{batch, 128, 28, 28, 128, 3, 1});

  53.     args.emplace_back(BenchArgs{batch, 128, 28, 28, 512, 1, 1});

  54. 

  55.     args.emplace_back(BenchArgs{batch, 512, 28, 28, 1024, 1, 2});

  56.     args.emplace_back(BenchArgs{batch, 512, 28, 28, 256, 1, 2});

  57.     args.emplace_back(BenchArgs{batch, 1024, 14, 14, 256, 1, 1});

  58.     args.emplace_back(BenchArgs{batch, 256, 14, 14, 256, 3, 1});

  59.     args.emplace_back(BenchArgs{batch, 256, 14, 14, 1024, 1, 1});

  60. 

  61.     args.emplace_back(BenchArgs{batch, 1024, 14, 14, 2048, 1, 2});

  62.     args.emplace_back(BenchArgs{batch, 1024, 14, 14, 512, 1, 2});

  63.     args.emplace_back(BenchArgs{batch, 2048, 7, 7, 512, 1, 1});

  64.     args.emplace_back(BenchArgs{batch, 512, 7, 7, 512, 3, 1});

  65.     args.emplace_back(BenchArgs{batch, 512, 7, 7, 2048, 1, 1});

  66.     return args;

  67. }

  68. 

  69. std::vector<BenchArgs> get_detection_bench_args(size_t batch = 16) {

  70.     std::vector<BenchArgs> args;

  71.     args.emplace_back(BenchArgs{batch, 4, 736, 1280, 8, 3, 2});

  72.     args.emplace_back(BenchArgs{batch, 32, 184, 320, 16, 3, 1});

  73.     args.emplace_back(BenchArgs{batch, 16, 184, 320, 32, 3, 1});

  74.     args.emplace_back(BenchArgs{batch, 8, 184, 320, 16, 3, 1});

  75.     args.emplace_back(BenchArgs{batch, 8, 184, 320, 32, 3, 1});

  76.     args.emplace_back(BenchArgs{batch, 64, 92, 160, 32, 3, 1});

  77.     args.emplace_back(BenchArgs{batch, 32, 184, 320, 64, 3, 2});

  78.     args.emplace_back(BenchArgs{batch, 32, 184, 320, 32, 3, 2});

  79.     args.emplace_back(BenchArgs{batch, 32, 92, 160, 64, 3, 1});

  80.     args.emplace_back(BenchArgs{batch, 64, 92, 160, 8, 3, 1});

  81.     args.emplace_back(BenchArgs{batch, 64, 92, 160, 128, 3, 2});

  82.     args.emplace_back(BenchArgs{batch, 128, 46, 80, 32, 3, 1});

  83.     args.emplace_back(BenchArgs{batch, 128, 46, 80, 256, 3, 2});

  84.     args.emplace_back(BenchArgs{batch, 128, 46, 80, 8, 3, 1});

  85.     args.emplace_back(BenchArgs{batch, 64, 92, 160, 32, 3, 2});

  86.     args.emplace_back(BenchArgs{batch, 32, 46, 80, 128, 3, 1});

  87.     args.emplace_back(BenchArgs{batch, 8, 46, 80, 32, 3, 1});

  88.     args.emplace_back(BenchArgs{batch, 64, 23, 40, 256, 3, 1});

  89.     args.emplace_back(BenchArgs{batch, 256, 23, 40, 64, 3, 1});

  90.     args.emplace_back(BenchArgs{batch, 128, 46, 80, 64, 3, 2});

  91.     args.emplace_back(BenchArgs{batch, 256, 23, 40, 8, 3, 1});

  92.     args.emplace_back(BenchArgs{batch, 8, 23, 40, 32, 3, 2});

  93.     args.emplace_back(BenchArgs{batch, 8, 12, 20, 8, 3, 1});

  94.     args.emplace_back(BenchArgs{batch, 8, 12, 20, 8, 3, 2});

  95.     args.emplace_back(BenchArgs{batch, 8, 6, 10, 8, 3, 1});

  96.     return args;

  97. }

  98. 

  99. void benchmark_target_algo(

  100.         Handle* handle, const std::vector<BenchArgs>& args, DType src_dtype,

  101.         DType filter_dtype, DType bias_dtype, DType dst_dtype,

  102.         const char* algo = nullptr,

  103.         param::ConvBias::Format format = param::ConvBias::Format::NCHW4) {

  104.     megdnn_assert(src_dtype.enumv() == filter_dtype.enumv());

  105.     CUBenchmarker<ConvBiasForward> benchmarker(handle);

  106.     CUBenchmarker<ConvBiasForward> benchmarker_cudnn(handle);

  107.     size_t RUNS = 1000;

  108.     benchmarker.set_display(false).set_times(RUNS);

  109.     benchmarker_cudnn.set_display(false).set_times(RUNS);

  110. 

  111. #define CUDNN_VERSION_STRING \

  112.     "v" V(CUDNN_MAJOR) "." V(CUDNN_MINOR) "." V(CUDNN_PATCHLEVEL)

  113.     benchmarker_cudnn.set_before_exec_callback(

  114.             conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  115.                     "DEFAULT:CUDNN:ConvBiasActivation:CUDNN_CONVOLUTION_FWD_"

  116.                     "ALGO_IMPLICIT_PRECOMP_"

  117.                     "GEMM" CUDNN_VERSION_STRING));

  118. 

  119.     benchmarker.set_dtype(0, src_dtype)

  120.             .set_dtype(1, filter_dtype)

  121.             .set_dtype(2, bias_dtype)

  122.             .set_dtype(3, dst_dtype)

  123.             .set_dtype(4, dst_dtype);

  124.     benchmarker_cudnn.set_dtype(0, src_dtype)

  125.             .set_dtype(1, filter_dtype)

  126.             .set_dtype(2, bias_dtype)

  127.             .set_dtype(3, dst_dtype)

  128.             .set_dtype(4, dst_dtype);

  129. 

  130.     using Param = ConvBias::Param;

  131.     using Format = Param::Format;

  132.     // helper function to change format

  133.     auto get_tensor_shape = [](TensorShape shape,

  134.                                Format format) -> TensorShape {

  135.         TensorShape ret;

  136.         if (format == Format::NCHW4) {

  137.             ret = static_cast<TensorShape>(

  138.                     TensorLayout{shape, dtype::Int8()}

  139.                             .reshape({shape[0], shape[1] / 4, 4, shape[2],

  140.                                       shape[3]})

  141.                             .dimshuffle({0, 1, 3, 4, 2}));

  142.         } else if (format == Format::CHWN4) {

  143.             ret = static_cast<TensorShape>(

  144.                     TensorLayout{shape, dtype::Int8()}

  145.                             .reshape({shape[0], shape[1] / 4, 4, shape[2],

  146.                                       shape[3]})

  147.                             .dimshuffle({1, 3, 4, 0, 2}));

  148.         }

  149.         return ret;

  150.     };

  151. 

  152.     for (auto&& arg : args) {

  153.         Param param;

  154.         param.pad_h = param.pad_w = arg.f / 2;

  155.         param.stride_h = param.stride_w = arg.s;

  156.         param.format = format;

  157. 

  158.         size_t ho = infer_conv_shape(arg.hi, arg.f, arg.s, arg.f / 2);

  159.         size_t wo = infer_conv_shape(arg.wi, arg.f, arg.s, arg.f / 2);

  160. 

  161.         benchmarker.set_param(param);

  162.         if (!algo) {

  163.             benchmarker.proxy()->target_algo = nullptr;

  164.         }

  165.         TensorShape src{arg.n, arg.ci, arg.hi, arg.wi},

  166.                 filter{arg.co, arg.ci, arg.f, arg.f}, bias{1, arg.co, 1, 1},

  167.                 z{arg.n, arg.co, ho, wo}, dst = z;

  168.         float time_in_ms = 0.f;

  169.         if (algo) {

  170.             time_in_ms =

  171.                     algo_benchmark<ConvBiasForward, OprProxy<ConvBiasForward>,

  172.                                    CUTimer>(benchmarker,

  173.                                             {get_tensor_shape(src, format),

  174.                                              get_tensor_shape(filter, format),

  175.                                              get_tensor_shape(bias, format),

  176.                                              {},

  177.                                              {}},

  178.                                             algo) /

  179.                     RUNS;

  180.         } else {

  181.             time_in_ms = benchmarker.execs({get_tensor_shape(src, format),

  182.                                             get_tensor_shape(filter, format),

  183.                                             get_tensor_shape(bias, format),

  184.                                             {},

  185.                                             {}}) /

  186.                          RUNS;

  187.         }

  188.         Format format_cudnn = Format::NCHW4;

  189.         param.format = format_cudnn;

  190.         benchmarker_cudnn.set_param(param);

  191.         auto time_in_ms_cudnn =

  192.                 benchmarker_cudnn.execs({get_tensor_shape(src, format_cudnn),

  193.                                          get_tensor_shape(filter, format_cudnn),

  194.                                          get_tensor_shape(bias, format_cudnn),

  195.                                          {},

  196.                                          {}}) /

  197.                 RUNS;

  198.         float flo = 2.0 * arg.n * arg.co * ho * wo * arg.ci * arg.f * arg.f /

  199.                     (1e12);

  200.         printf("src=%s, filter=%s, dst=%s, time(algo=%s)=%.2f %.2fTops, "

  201.                "time(cudnn)=%.2f %.2fTops, "

  202.                "perf(algo=%s)/perf(cudnn)=%.2f\n",

  203.                src.to_string().c_str(), filter.to_string().c_str(),

  204.                dst.to_string().c_str(), algo, time_in_ms,

  205.                (flo / (time_in_ms * 1e-3)), time_in_ms_cudnn,

  206.                (flo / (time_in_ms_cudnn * 1e-3)), algo,

  207.                time_in_ms_cudnn / time_in_ms);

  208.         printf("bench with z tensor\n");

  209.         if (algo) {

  210.             time_in_ms =

  211.                     algo_benchmark<ConvBiasForward, OprProxy<ConvBiasForward>,

  212.                                    CUTimer>(benchmarker,

  213.                                             {get_tensor_shape(src, format),

  214.                                              get_tensor_shape(filter, format),

  215.                                              get_tensor_shape(bias, format),

  216.                                              get_tensor_shape(z, format),

  217.                                              {}},

  218.                                             algo) /

  219.                     RUNS;

  220.         } else {

  221.             time_in_ms = benchmarker.execs({get_tensor_shape(src, format),

  222.                                             get_tensor_shape(filter, format),

  223.                                             get_tensor_shape(bias, format),

  224.                                             get_tensor_shape(z, format),

  225.                                             {}}) /

  226.                          RUNS;

  227.         }

  228.         time_in_ms_cudnn =

  229.                 benchmarker_cudnn.execs({get_tensor_shape(src, format_cudnn),

  230.                                          get_tensor_shape(filter, format_cudnn),

  231.                                          get_tensor_shape(bias, format_cudnn),

  232.                                          get_tensor_shape(z, format_cudnn),

  233.                                          {}}) /

  234.                 RUNS;

  235.         printf("src=%s, filter=%s, dst=%s, time(algo=%s)=%.2f %.2fTops, "

  236.                "time(cudnn)=%.2f %.2fTops, "

  237.                "perf(algo=%s)/perf(cudnn)=%.2f\n",

  238.                src.to_string().c_str(), filter.to_string().c_str(),

  239.                dst.to_string().c_str(), algo, time_in_ms,

  240.                (flo / (time_in_ms * 1e-3)), time_in_ms_cudnn,

  241.                (flo / (time_in_ms_cudnn * 1e-3)), algo,

  242.                time_in_ms_cudnn / time_in_ms);

  243.     }

  244. }

  245. 

  246. void benchmark_target_algo_with_cudnn_tsc(

  247.         Handle* handle, const std::vector<BenchArgs>& args, DType src_dtype,

  248.         DType filter_dtype, DType bias_dtype, DType dst_dtype,

  249.         const char* algo = nullptr,

  250.         param::ConvBias::Format format = param::ConvBias::Format::NCHW4) {

  251.     megdnn_assert(src_dtype.enumv() == filter_dtype.enumv());

  252.     CUBenchmarker<ConvBiasForward> benchmarker(handle);

  253.     CUBenchmarker<ConvBiasForward> benchmarker_cudnn(handle);

  254.     size_t RUNS = 1000;

  255.     benchmarker.set_display(false).set_times(RUNS);

  256.     benchmarker_cudnn.set_display(false).set_times(RUNS);

  257. 

  258.     std::unique_ptr<OprProxy<ConvBiasForward>> proxy{

  259.             new OprProxy<ConvBiasForward>{true}};

  260. 

  261.     if (!algo) {

  262.         benchmarker.set_proxy(proxy);

  263.     }

  264. 

  265.     benchmarker_cudnn.set_before_exec_callback(

  266.             conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  267.                     "DEFAULT:CUDNN:ConvBiasActivation:CUDNN_CONVOLUTION_FWD_"

  268.                     "ALGO_IMPLICIT_PRECOMP_"

  269.                     "GEMM" CUDNN_VERSION_STRING));

  270. #undef CUDNN_VERSION_STRING

  271. 

  272.     benchmarker.set_dtype(0, src_dtype)

  273.             .set_dtype(1, filter_dtype)

  274.             .set_dtype(2, bias_dtype)

  275.             .set_dtype(3, dst_dtype)

  276.             .set_dtype(4, dst_dtype);

  277.     benchmarker_cudnn.set_dtype(0, src_dtype)

  278.             .set_dtype(1, filter_dtype)

  279.             .set_dtype(2, bias_dtype)

  280.             .set_dtype(3, dst_dtype)

  281.             .set_dtype(4, dst_dtype);

  282. 

  283.     using Param = ConvBias::Param;

  284.     using Format = Param::Format;

  285.     // helper function to change format

  286.     auto get_tensor_shape = [](TensorShape shape,

  287.                                Format format) -> TensorShape {

  288.         TensorShape ret;

  289.         if (format == Format::NCHW4) {

  290.             ret = static_cast<TensorShape>(

  291.                     TensorLayout{shape, dtype::Int8()}

  292.                             .reshape({shape[0], shape[1] / 4, 4, shape[2],

  293.                                       shape[3]})

  294.                             .dimshuffle({0, 1, 3, 4, 2}));

  295.         } else if (format == Format::NCHW32) {

  296.             ret = static_cast<TensorShape>(

  297.                     TensorLayout{shape, dtype::Int8()}

  298.                             .reshape({shape[0], shape[1] / 32, 32, shape[2],

  299.                                       shape[3]})

  300.                             .dimshuffle({0, 1, 3, 4, 2}));

  301.         } else if (format == Format::CHWN4) {

  302.             ret = static_cast<TensorShape>(

  303.                     TensorLayout{shape, dtype::Int8()}

  304.                             .reshape({shape[0], shape[1] / 4, 4, shape[2],

  305.                                       shape[3]})

  306.                             .dimshuffle({1, 3, 4, 0, 2}));

  307.         }

  308.         return ret;

  309.     };

  310. 

  311.     for (auto&& arg : args) {

  312.         Param param;

  313.         param.pad_h = param.pad_w = arg.f / 2;

  314.         param.stride_h = param.stride_w = arg.s;

  315.         param.format = format;

  316. 

  317.         size_t ho = infer_conv_shape(arg.hi, arg.f, arg.s, arg.f / 2);

  318.         size_t wo = infer_conv_shape(arg.wi, arg.f, arg.s, arg.f / 2);

  319. 

  320.         benchmarker.set_param(param);

  321.         if (!algo) {

  322.             benchmarker.proxy()->target_algo = nullptr;

  323.         }

  324.         TensorShape src{arg.n, arg.ci, arg.hi, arg.wi},

  325.                 filter{arg.co, arg.ci, arg.f, arg.f}, bias{1, arg.co, 1, 1},

  326.                 z{arg.n, arg.co, ho, wo}, dst = z;

  327.         // skip testcase which cannot enable nchw32 tensorcore

  328.         if (format == Format::NCHW32 && (arg.co % 32 != 0 || arg.ci % 32 != 0))

  329.             continue;

  330.         // skip testcase which cannot enable nchw4/chwn4 tensorcore

  331.         if ((format == Format::CHWN4 || format == Format::NCHW4) &&

  332.             (arg.ci % 16 != 0))

  333.             continue;

  334.         float time_in_ms = 0.f;

  335.         if (algo) {

  336.             time_in_ms =

  337.                     algo_benchmark<ConvBiasForward, OprProxy<ConvBiasForward>,

  338.                                    CUTimer>(benchmarker,

  339.                                             {get_tensor_shape(src, format),

  340.                                              get_tensor_shape(filter, format),

  341.                                              get_tensor_shape(bias, format),

  342.                                              {},

  343.                                              {}},

  344.                                             algo) /

  345.                     RUNS;

  346.         } else {

  347.             time_in_ms = benchmarker.execs({get_tensor_shape(src, format),

  348.                                             get_tensor_shape(filter, format),

  349.                                             get_tensor_shape(bias, format),

  350.                                             {},

  351.                                             {}}) /

  352.                          RUNS;

  353.         }

  354.         Format format_cudnn = arg.ci % 32 == 0 && arg.co % 32 == 0

  355.                                       ? Format::NCHW32

  356.                                       : Format::NCHW4;

  357.         param.format = format_cudnn;

  358.         benchmarker_cudnn.set_param(param);

  359.         auto time_in_ms_cudnn =

  360.                 benchmarker_cudnn.execs({get_tensor_shape(src, format_cudnn),

  361.                                          get_tensor_shape(filter, format_cudnn),

  362.                                          get_tensor_shape(bias, format_cudnn),

  363.                                          {},

  364.                                          {}}) /

  365.                 RUNS;

  366.         float flo = 2.0 * arg.n * arg.co * ho * wo * arg.ci * arg.f * arg.f /

  367.                     (1e12);

  368.         printf("src=%s, filter=%s, dst=%s, time(algo=%s)=%.2f %.2fTops, "

  369.                "time(cudnn)=%.2f %.2fTops, "

  370.                "perf(algo=%s)/perf(cudnn)=%.2f\n",

  371.                src.to_string().c_str(), filter.to_string().c_str(),

  372.                dst.to_string().c_str(), algo, time_in_ms,

  373.                (flo / (time_in_ms * 1e-3)), time_in_ms_cudnn,

  374.                (flo / (time_in_ms_cudnn * 1e-3)), algo,

  375.                time_in_ms_cudnn / time_in_ms);

  376.         printf("bench with z tensor\n");

  377.         if (algo) {

  378.             time_in_ms =

  379.                     algo_benchmark<ConvBiasForward, OprProxy<ConvBiasForward>,

  380.                                    CUTimer>(benchmarker,

  381.                                             {get_tensor_shape(src, format),

  382.                                              get_tensor_shape(filter, format),

  383.                                              get_tensor_shape(bias, format),

  384.                                              get_tensor_shape(z, format),

  385.                                              {}},

  386.                                             algo) /

  387.                     RUNS;

  388.         } else {

  389.             time_in_ms = benchmarker.execs({get_tensor_shape(src, format),

  390.                                             get_tensor_shape(filter, format),

  391.                                             get_tensor_shape(bias, format),

  392.                                             get_tensor_shape(z, format),

  393.                                             {}}) /

  394.                          RUNS;

  395.         }

  396.         time_in_ms_cudnn =

  397.                 benchmarker_cudnn.execs({get_tensor_shape(src, format_cudnn),

  398.                                          get_tensor_shape(filter, format_cudnn),

  399.                                          get_tensor_shape(bias, format_cudnn),

  400.                                          get_tensor_shape(z, format_cudnn),

  401.                                          {}}) /

  402.                 RUNS;

  403.         printf("src=%s, filter=%s, dst=%s, time(algo=%s)=%.2f %.2fTops, "

  404.                "time(cudnn)=%.2f %.2fTops, "

  405.                "perf(algo=%s)/perf(cudnn)=%.2f\n",

  406.                src.to_string().c_str(), filter.to_string().c_str(),

  407.                dst.to_string().c_str(), algo, time_in_ms,

  408.                (flo / (time_in_ms * 1e-3)), time_in_ms_cudnn,

  409.                (flo / (time_in_ms_cudnn * 1e-3)), algo,

  410.                time_in_ms_cudnn / time_in_ms);

  411.     }

  412. }

  413. #endif

  414. }  // namespace

  415. 

  416. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_1x1) {

  417.     require_compute_capability(6, 1);

  418.     conv_bias::check_conv_bias(

  419.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  420.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  421.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

  422.             param::ConvBias::Format::NCHW4, conv_bias::get_int8_nchw4_args(1));

  423. }

  424. 

  425. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_3x3) {

  426.     require_compute_capability(6, 1);

  427.     conv_bias::check_conv_bias(

  428.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  429.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  430.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

  431.             param::ConvBias::Format::NCHW4);

  432. }

  433. 

  434. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_5x5) {

  435.     require_compute_capability(6, 1);

  436.     conv_bias::check_conv_bias(

  437.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  438.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  439.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

  440.             param::ConvBias::Format::NCHW4, conv_bias::get_int8_nchw4_args(5));

  441. }

  442. 

  443. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_7x7) {

  444.     require_compute_capability(6, 1);

  445.     conv_bias::check_conv_bias(

  446.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  447.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  448.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

  449.             param::ConvBias::Format::NCHW4, conv_bias::get_int8_nchw4_args(7));

  450. }

  451. 

  452. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_WITH_Z) {

  453.     require_compute_capability(6, 1);

  454.     Checker<ConvBiasForward> checker(handle_cuda());

  455.     checker.set_before_exec_callback(

  456.             conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  457.                     "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM"));

  458.     UniformIntRNG rng{-3, 3};

  459.     UniformIntRNG bias_rng{-50, 50};

  460.     checker.set_rng(0, &rng)

  461.             .set_rng(1, &rng)

  462.             .set_rng(2, &bias_rng)

  463.             .set_rng(3, &rng)

  464.             .set_dtype(0, dtype::QuantizedS8{1.2f})

  465.             .set_dtype(1, dtype::QuantizedS8{1.3f})

  466.             .set_dtype(2, dtype::QuantizedS32{1.2f * 1.3f})

  467.             .set_dtype(3, dtype::QuantizedS8{1.1f})

  468.             .set_dtype(4, dtype::QuantizedS8{1.0f})

  469.             .set_epsilon(1 + 1e-3)

  470.             .set_max_avg_error(1e-1)

  471.             .set_max_avg_biased_error(1e-1);

  472.     param::ConvBias param;

  473.     param.pad_h = param.pad_w = 1;

  474.     param.stride_h = param.stride_w = 1;

  475.     param.format = param::ConvBias::Format::NCHW4;

  476.     checker.set_param(param).execs({{32, 4, 12, 12, 4},

  477.                                     {16, 4, 3, 3, 4},

  478.                                     {1, 4, 1, 1, 4},

  479.                                     {32, 4, 12, 12, 4},

  480.                                     {}});

  481. }

  482. 

  483. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_STRIDE2_WITH_Z) {

  484.     require_compute_capability(6, 1);

  485.     Checker<ConvBiasForward> checker(handle_cuda());

  486.     checker.set_before_exec_callback(

  487.             conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  488.                     "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM"));

  489.     UniformIntRNG rng{-3, 3};

  490.     UniformIntRNG bias_rng{-50, 50};

  491.     checker.set_rng(0, &rng)

  492.             .set_rng(1, &rng)

  493.             .set_rng(2, &bias_rng)

  494.             .set_rng(3, &rng)

  495.             .set_dtype(0, dtype::QuantizedS8{1.2f})

  496.             .set_dtype(1, dtype::QuantizedS8{1.3f})

  497.             .set_dtype(2, dtype::QuantizedS32{1.2f * 1.3f})

  498.             .set_dtype(3, dtype::QuantizedS8{1.1f})

  499.             .set_dtype(4, dtype::QuantizedS8{1.0f})

  500.             .set_epsilon(1 + 1e-3)

  501.             .set_max_avg_error(1e-1)

  502.             .set_max_avg_biased_error(1e-1);

  503.     param::ConvBias param;

  504.     param.pad_h = param.pad_w = 1;

  505.     param.stride_h = param.stride_w = 2;

  506.     param.format = param::ConvBias::Format::NCHW4;

  507.     checker.set_param(param).execs({{32, 4, 12, 12, 4},

  508.                                     {16, 4, 3, 3, 4},

  509.                                     {1, 4, 1, 1, 4},

  510.                                     {32, 4, 6, 6, 4},

  511.                                     {}});

  512. }

  513. 

  514. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_CHECK_BOUNDS_1x1) {

  515.     require_compute_capability(6, 1);

  516.     conv_bias::check_conv_bias(

  517.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  518.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  519.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

  520.             param::ConvBias::Format::NCHW4,

  521.             conv_bias::get_int8_nchw4_args_check_bounds(1));

  522. }

  523. 

  524. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_CHECK_BOUNDS_3x3) {

  525.     require_compute_capability(6, 1);

  526.     conv_bias::check_conv_bias(

  527.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  528.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  529.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

  530.             param::ConvBias::Format::NCHW4,

  531.             conv_bias::get_int8_nchw4_args_check_bounds(3));

  532. }

  533. 

  534. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_CHECK_BOUNDS_5x5) {

  535.     require_compute_capability(6, 1);

  536.     conv_bias::check_conv_bias(

  537.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  538.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  539.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

  540.             param::ConvBias::Format::NCHW4,

  541.             conv_bias::get_int8_nchw4_args_check_bounds(5));

  542. }

  543. 

  544. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_CHECK_BOUNDS_7x7) {

  545.     require_compute_capability(6, 1);

  546.     conv_bias::check_conv_bias(

  547.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  548.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  549.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

  550.             param::ConvBias::Format::NCHW4,

  551.             conv_bias::get_int8_nchw4_args_check_bounds(7));

  552. }

  553. 

  554. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4) {

  555.     require_compute_capability(6, 1);

  556.     conv_bias::check_conv_bias(

  557.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  558.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  559.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  560.             param::ConvBias::Format::CHWN4);

  561. }

  562. 

  563. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_WITH_Z) {

  564.     require_compute_capability(6, 1);

  565.     Checker<ConvBiasForward> checker(handle_cuda());

  566.     checker.set_before_exec_callback(

  567.             conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  568.                     "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM"));

  569.     UniformIntRNG rng{-3, 3};

  570.     UniformIntRNG bias_rng{-50, 50};

  571.     checker.set_rng(0, &rng)

  572.             .set_rng(1, &rng)

  573.             .set_rng(2, &bias_rng)

  574.             .set_rng(3, &rng)

  575.             .set_dtype(0, dtype::QuantizedS8{1.2f})

  576.             .set_dtype(1, dtype::QuantizedS8{1.3f})

  577.             .set_dtype(2, dtype::QuantizedS32{1.2f * 1.3f})

  578.             .set_dtype(3, dtype::QuantizedS8{1.1f})

  579.             .set_dtype(4, dtype::QuantizedS8{1.1f})

  580.             .set_epsilon(1 + 1e-3)

  581.             .set_max_avg_error(1e-1)

  582.             .set_max_avg_biased_error(1e-1);

  583.     param::ConvBias param;

  584.     param.pad_h = param.pad_w = 1;

  585.     param.stride_h = param.stride_w = 1;

  586.     param.format = param::ConvBias::Format::CHWN4;

  587.     checker.set_param(param).execs({{4, 12, 12, 32, 4},

  588.                                     {4, 3, 3, 16, 4},

  589.                                     {4, 1, 1, 1, 4},

  590.                                     {4, 12, 12, 32, 4},

  591.                                     {}});

  592. }

  593. 

  594. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_HSWISH) {

  595.     require_compute_capability(6, 1);

  596.     Checker<ConvBiasForward> checker(handle_cuda());

  597.     checker.set_before_exec_callback(

  598.             conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  599.                     "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM"));

  600.     UniformIntRNG rng{-3, 3};

  601.     UniformIntRNG bias_rng{-50, 50};

  602.     checker.set_rng(0, &rng)

  603.             .set_rng(1, &rng)

  604.             .set_rng(2, &bias_rng)

  605.             .set_rng(3, &rng)

  606.             .set_dtype(0, dtype::QuantizedS8{1.2f})

  607.             .set_dtype(1, dtype::QuantizedS8{1.3f})

  608.             .set_dtype(2, dtype::QuantizedS32{1.2f * 1.3f})

  609.             .set_dtype(4, dtype::QuantizedS8{0.001f})

  610.             .set_epsilon(1 + 1e-3)

  611.             .set_max_avg_error(1e-1)

  612.             .set_max_avg_biased_error(1e-1);

  613.     param::ConvBias param;

  614.     param.pad_h = param.pad_w = 1;

  615.     param.stride_h = param.stride_w = 1;

  616.     param.format = param::ConvBias::Format::CHWN4;

  617.     param.nonlineMode = param::ConvBias::NonlineMode::H_SWISH;

  618.     checker.set_param(param).execs(

  619.             {{4, 12, 12, 32, 4}, {4, 3, 3, 16, 4}, {4, 1, 1, 1, 4}, {}, {}});

  620. }

  621. 

  622. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_CHECK_BOUNDS) {

  623.     require_compute_capability(6, 1);

  624.     conv_bias::check_conv_bias(

  625.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  626.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  627.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  628.             param::ConvBias::Format::CHWN4,

  629.             conv_bias::get_int8_chwn4_args_check_bounds(3));

  630. }

  631. 

  632. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_SMALL_CHANNEL_1x1) {

  633.     require_compute_capability(6, 1);

  634.     conv_bias::check_conv_bias(

  635.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  636.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  637.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  638.             param::ConvBias::Format::CHWN4,

  639.             conv_bias::get_int8_chwn4_small_channel_args(1));

  640. }

  641. 

  642. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_SMALL_CHANNEL_3x3) {

  643.     require_compute_capability(6, 1);

  644.     conv_bias::check_conv_bias(

  645.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  646.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  647.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  648.             param::ConvBias::Format::CHWN4,

  649.             conv_bias::get_int8_chwn4_small_channel_args(3));

  650. }

  651. 

  652. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_SMALL_CHANNEL_5x5) {

  653.     require_compute_capability(6, 1);

  654.     conv_bias::check_conv_bias(

  655.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  656.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  657.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  658.             param::ConvBias::Format::CHWN4,

  659.             conv_bias::get_int8_chwn4_small_channel_args(5));

  660. }

  661. 

  662. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_SMALL_CHANNEL_7x7) {

  663.     require_compute_capability(6, 1);

  664.     conv_bias::check_conv_bias(

  665.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  666.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  667.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  668.             param::ConvBias::Format::CHWN4,

  669.             conv_bias::get_int8_chwn4_small_channel_args(7));

  670. }

  671. 

  672. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_SMALL_CHANNEL_CHECK_BOUNDS) {

  673.     require_compute_capability(6, 1);

  674.     conv_bias::check_conv_bias(

  675.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  676.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  677.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

  678.             param::ConvBias::Format::NCHW4,

  679.             conv_bias::get_int8_nchw4_small_channel_args_check_bounds(3));

  680. }

  681. 

  682. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_SMALL_CHANNEL_1x1_CHECK_BOUNDS) {

  683.     require_compute_capability(6, 1);

  684.     conv_bias::check_conv_bias(

  685.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  686.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  687.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  688.             param::ConvBias::Format::CHWN4,

  689.             conv_bias::get_int8_chwn4_small_channel_args_check_bounds(1));

  690. }

  691. 

  692. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_SMALL_CHANNEL_5x5_CHECK_BOUNDS) {

  693.     require_compute_capability(6, 1);

  694.     conv_bias::check_conv_bias(

  695.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  696.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  697.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  698.             param::ConvBias::Format::CHWN4,

  699.             conv_bias::get_int8_chwn4_small_channel_args_check_bounds(5));

  700. }

  701. 

  702. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_SMALL_CHANNEL_7x7_CHECK_BOUNDS) {

  703.     require_compute_capability(6, 1);

  704.     conv_bias::check_conv_bias(

  705.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  706.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  707.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  708.             param::ConvBias::Format::CHWN4,

  709.             conv_bias::get_int8_chwn4_small_channel_args_check_bounds(7));

  710. }

  711. 

  712. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_TENSORCORE_1x1) {

  713.     require_compute_capability(7, 5);

  714.     conv_bias::check_conv_bias(

  715.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  716.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  717.             handle_cuda(), "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  718.             param::ConvBias::Format::NCHW4,

  719.             conv_bias::get_int8_nchw4_tensorcore_args(1));

  720. }

  721. 

  722. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_TENSORCORE_3x3) {

  723.     require_compute_capability(7, 5);

  724.     conv_bias::check_conv_bias(

  725.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  726.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  727.             handle_cuda(), "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  728.             param::ConvBias::Format::NCHW4,

  729.             conv_bias::get_int8_nchw4_tensorcore_args(3));

  730. }

  731. 

  732. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_TENSORCORE_5x5) {

  733.     require_compute_capability(7, 5);

  734.     conv_bias::check_conv_bias(

  735.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  736.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  737.             handle_cuda(), "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  738.             param::ConvBias::Format::NCHW4,

  739.             conv_bias::get_int8_nchw4_tensorcore_args(5));

  740. }

  741. 

  742. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_TENSORCORE_7x7) {

  743.     require_compute_capability(7, 5);

  744.     conv_bias::check_conv_bias(

  745.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  746.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  747.             handle_cuda(), "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  748.             param::ConvBias::Format::NCHW4,

  749.             conv_bias::get_int8_nchw4_tensorcore_args(7));

  750. }

  751. 

  752. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_TENSORCORE_CHECK_BOUNDS_ALGO_0) {

  753.     require_compute_capability(7, 5);

  754.     conv_bias::check_conv_bias(

  755.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  756.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  757.             handle_cuda(), "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  758.             param::ConvBias::Format::NCHW4,

  759.             conv_bias::get_int8_nchw4_args_check_bounds(3));

  760. }

  761. 

  762. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_TENSORCORE_CHECK_BOUNDS_ALGO_1) {

  763.     require_compute_capability(7, 5);

  764.     conv_bias::check_conv_bias(

  765.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  766.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  767.             handle_cuda(), "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma8x32x16",

  768.             param::ConvBias::Format::NCHW4,

  769.             conv_bias::get_int8_nchw4_args_check_bounds(3));

  770. }

  771. 

  772. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_TENSORCORE_CHECK_BOUNDS_ALGO_2) {

  773.     require_compute_capability(7, 5);

  774.     conv_bias::check_conv_bias(

  775.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  776.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  777.             handle_cuda(), "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma32x8x16",

  778.             param::ConvBias::Format::NCHW4,

  779.             conv_bias::get_int8_nchw4_args_check_bounds(3));

  780. }

  781. 

  782. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_TENSORCORE_ALGO_0) {

  783.     require_compute_capability(7, 5);

  784.     conv_bias::check_conv_bias(

  785.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  786.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.1f},

  787.             handle_cuda(), "INT8_CHWN4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  788.             param::ConvBias::Format::CHWN4,

  789.             conv_bias::get_int8_chwn4_tensorcore_args(3));

  790. }

  791. 

  792. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_TENSORCORE_ALGO_1) {

  793.     require_compute_capability(7, 5);

  794.     conv_bias::check_conv_bias(

  795.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  796.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.1f},

  797.             handle_cuda(), "INT8_CHWN4_IMMA_IMPLICIT_GEMM_mma32x8x16",

  798.             param::ConvBias::Format::CHWN4,

  799.             conv_bias::get_int8_chwn4_tensorcore_args(3));

  800. }

  801. 

  802. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_TENSORCORE_ALGO_2) {

  803.     require_compute_capability(7, 5);

  804.     conv_bias::check_conv_bias(

  805.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  806.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.1f},

  807.             handle_cuda(), "INT8_CHWN4_IMMA_IMPLICIT_GEMM_mma8x32x16",

  808.             param::ConvBias::Format::CHWN4,

  809.             conv_bias::get_int8_chwn4_tensorcore_args(3));

  810. }

  811. 

  812. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_TENSORCORE_CHECK_BOUNDS_1x1) {

  813.     require_compute_capability(7, 5);

  814.     conv_bias::check_conv_bias(

  815.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  816.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  817.             handle_cuda(), "INT8_CHWN4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  818.             param::ConvBias::Format::CHWN4,

  819.             conv_bias::get_int8_chwn4_args_check_bounds(1));

  820. }

  821. 

  822. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_TENSORCORE_CHECK_BOUNDS_5x5) {

  823.     require_compute_capability(7, 5);

  824.     conv_bias::check_conv_bias(

  825.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  826.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  827.             handle_cuda(), "INT8_CHWN4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  828.             param::ConvBias::Format::CHWN4,

  829.             conv_bias::get_int8_chwn4_args_check_bounds(5));

  830. }

  831. 

  832. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_TENSORCORE_CHECK_BOUNDS_7x7) {

  833.     require_compute_capability(7, 5);

  834.     conv_bias::check_conv_bias(

  835.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  836.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  837.             handle_cuda(), "INT8_CHWN4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  838.             param::ConvBias::Format::CHWN4,

  839.             conv_bias::get_int8_chwn4_args_check_bounds(7));

  840. }

  841. 

  842. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_TENSORCORE_WITH_Z) {

  843.     require_compute_capability(7, 5);

  844.     Checker<ConvBiasForward> checker(handle_cuda());

  845.     checker.set_before_exec_callback(

  846.             conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  847.                     "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma16x16x16"));

  848.     UniformIntRNG rng{-3, 3};

  849.     UniformIntRNG bias_rng{-50, 50};

  850.     checker.set_rng(0, &rng)

  851.             .set_rng(1, &rng)

  852.             .set_rng(2, &bias_rng)

  853.             .set_rng(3, &rng)

  854.             .set_dtype(0, dtype::QuantizedS8{1.2f})

  855.             .set_dtype(1, dtype::QuantizedS8{1.3f})

  856.             .set_dtype(2, dtype::QuantizedS32{1.2f * 1.3f})

  857.             .set_dtype(3, dtype::QuantizedS8{1.1f})

  858.             .set_dtype(4, dtype::QuantizedS8{1.0f})

  859.             .set_epsilon(1 + 1e-3)

  860.             .set_max_avg_error(1e-1)

  861.             .set_max_avg_biased_error(1e-1);

  862.     param::ConvBias param;

  863.     param.pad_h = param.pad_w = 1;

  864.     param.stride_h = param.stride_w = 1;

  865.     param.format = param::ConvBias::Format::NCHW4;

  866.     checker.set_param(param).execs({{64, 8, 12, 12, 4},

  867.                                     {64, 8, 3, 3, 4},

  868.                                     {1, 16, 1, 1, 4},

  869.                                     {64, 16, 12, 12, 4},

  870.                                     {}});

  871. }

  872. 

  873. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_TENSORCORE_WITH_Z) {

  874.     require_compute_capability(7, 5);

  875.     Checker<ConvBiasForward> checker(handle_cuda());

  876.     checker.set_before_exec_callback(

  877.             conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  878.                     "INT8_CHWN4_IMMA_IMPLICIT_GEMM_mma16x16x16"));

  879.     UniformIntRNG rng{-3, 3};

  880.     UniformIntRNG bias_rng{-50, 50};

  881.     checker.set_rng(0, &rng)

  882.             .set_rng(1, &rng)

  883.             .set_rng(2, &bias_rng)

  884.             .set_rng(3, &rng)

  885.             .set_dtype(0, dtype::QuantizedS8{1.2f})

  886.             .set_dtype(1, dtype::QuantizedS8{1.3f})

  887.             .set_dtype(2, dtype::QuantizedS32{1.2f * 1.3f})

  888.             .set_dtype(3, dtype::QuantizedS8{1.1f})

  889.             .set_dtype(4, dtype::QuantizedS8{1.0f})

  890.             .set_epsilon(1 + 1e-3)

  891.             .set_max_avg_error(1e-1)

  892.             .set_max_avg_biased_error(1e-1);

  893.     param::ConvBias param;

  894.     param.pad_h = param.pad_w = 1;

  895.     param.stride_h = param.stride_w = 1;

  896.     param.format = param::ConvBias::Format::CHWN4;

  897.     checker.set_param(param).execs({{8, 12, 12, 64, 4},

  898.                                     {8, 3, 3, 64, 4},

  899.                                     {16, 1, 1, 1, 4},

  900.                                     {16, 12, 12, 64, 4},

  901.                                     {}});

  902. }

  903. 

  904. TEST_F(CUDA,

  905.        CONV_BIAS_INT8_CHWN4_REFORMAT_FILTER_TENSORCORE_CHECK_BOUNDS_ALGO_0) {

  906.     require_compute_capability(7, 5);

  907.     conv_bias::check_conv_bias(

  908.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  909.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  910.             handle_cuda(),

  911.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_REORDER_FILTER_mma16x16x16",

  912.             param::ConvBias::Format::CHWN4,

  913.             conv_bias::get_int8_chwn4_args_check_bounds(3));

  914. }

  915. 

  916. TEST_F(CUDA,

  917.        CONV_BIAS_INT8_CHWN4_REFORMAT_FILTER_TENSORCORE_CHECK_BOUNDS_ALGO_1) {

  918.     require_compute_capability(7, 5);

  919.     conv_bias::check_conv_bias(

  920.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  921.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  922.             handle_cuda(),

  923.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_REORDER_FILTER_mma8x32x16",

  924.             param::ConvBias::Format::CHWN4,

  925.             conv_bias::get_int8_chwn4_args_check_bounds(3));

  926. }

  927. 

  928. TEST_F(CUDA,

  929.        CONV_BIAS_INT8_CHWN4_REFORMAT_FILTER_TENSORCORE_CHECK_BOUNDS_ALGO_2) {

  930.     require_compute_capability(7, 5);

  931.     conv_bias::check_conv_bias(

  932.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  933.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  934.             handle_cuda(),

  935.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_REORDER_FILTER_mma32x8x16",

  936.             param::ConvBias::Format::CHWN4,

  937.             conv_bias::get_int8_chwn4_args_check_bounds(3));

  938. }

  939. 

  940. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_REFORMAT_FILTER_TENSORCORE_ALGO_0) {

  941.     require_compute_capability(7, 5);

  942.     conv_bias::check_conv_bias(

  943.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  944.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  945.             handle_cuda(),

  946.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_REORDER_FILTER_mma16x16x16",

  947.             param::ConvBias::Format::CHWN4, conv_bias::get_int8_chwn4_args(3));

  948. }

  949. 

  950. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_REFORMAT_FILTER_TENSORCORE_ALGO_1) {

  951.     require_compute_capability(7, 5);

  952.     conv_bias::check_conv_bias(

  953.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  954.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  955.             handle_cuda(),

  956.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_REORDER_FILTER_mma8x32x16",

  957.             param::ConvBias::Format::CHWN4, conv_bias::get_int8_chwn4_args(3));

  958. }

  959. 

  960. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_REFORMAT_FILTER_TENSORCORE_ALGO_2) {

  961.     require_compute_capability(7, 5);

  962.     conv_bias::check_conv_bias(

  963.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  964.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  965.             handle_cuda(),

  966.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_REORDER_FILTER_mma32x8x16",

  967.             param::ConvBias::Format::CHWN4, conv_bias::get_int8_chwn4_args(3));

  968. }

  969. 

  970. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_ALGO_0) {

  971.     require_compute_capability(7, 5);

  972.     conv_bias::check_conv_bias(

  973.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  974.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.1f},

  975.             handle_cuda(),

  976.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma16x16x16",

  977.             param::ConvBias::Format::CHWN4, conv_bias::get_int8_chwn4_args(3));

  978. }

  979. 

  980. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_ALGO_1) {

  981.     require_compute_capability(7, 5);

  982.     conv_bias::check_conv_bias(

  983.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  984.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  985.             handle_cuda(),

  986.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma8x32x16",

  987.             param::ConvBias::Format::CHWN4, conv_bias::get_int8_chwn4_args(3));

  988. }

  989. 

  990. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_ALGO_2) {

  991.     require_compute_capability(7, 5);

  992.     conv_bias::check_conv_bias(

  993.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  994.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.3f},

  995.             handle_cuda(),

  996.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma32x8x16",

  997.             param::ConvBias::Format::CHWN4, conv_bias::get_int8_chwn4_args(3));

  998. }

  999. 

  1000. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_1x1) {

  1001.     require_compute_capability(7, 5);

  1002.     conv_bias::check_conv_bias(

  1003.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  1004.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.1f},

  1005.             handle_cuda(),

  1006.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma16x16x16",

  1007.             param::ConvBias::Format::CHWN4, conv_bias::get_int8_chwn4_args(1));

  1008. }

  1009. 

  1010. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_5x5) {

  1011.     require_compute_capability(7, 5);

  1012.     conv_bias::check_conv_bias(

  1013.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  1014.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.1f},

  1015.             handle_cuda(),

  1016.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma16x16x16",

  1017.             param::ConvBias::Format::CHWN4,

  1018.             conv_bias::get_int8_chwn4_args_small_batch(5));

  1019. }

  1020. 

  1021. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_7x7) {

  1022.     require_compute_capability(7, 5);

  1023.     conv_bias::check_conv_bias(

  1024.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  1025.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.1f},

  1026.             handle_cuda(),

  1027.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma16x16x16",

  1028.             param::ConvBias::Format::CHWN4,

  1029.             conv_bias::get_int8_chwn4_args_small_batch(7));

  1030. }

  1031. 

  1032. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_5x5_ALGO_1) {

  1033.     require_compute_capability(7, 5);

  1034.     conv_bias::check_conv_bias(

  1035.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  1036.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.1f},

  1037.             handle_cuda(),

  1038.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma32x8x16",

  1039.             param::ConvBias::Format::CHWN4,

  1040.             conv_bias::get_int8_chwn4_args_small_batch(5));

  1041. }

  1042. 

  1043. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_5x5_ALGO_2) {

  1044.     require_compute_capability(7, 5);

  1045.     conv_bias::check_conv_bias(

  1046.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  1047.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.1f},

  1048.             handle_cuda(),

  1049.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma8x32x16",

  1050.             param::ConvBias::Format::CHWN4,

  1051.             conv_bias::get_int8_chwn4_args_small_batch(5));

  1052. }

  1053. 

  1054. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_1x1_ALGO_1) {

  1055.     require_compute_capability(7, 5);

  1056.     conv_bias::check_conv_bias(

  1057.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  1058.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.1f},

  1059.             handle_cuda(),

  1060.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma32x8x16",

  1061.             param::ConvBias::Format::CHWN4,

  1062.             conv_bias::get_int8_chwn4_args_small_batch(1));

  1063. }

  1064. 

  1065. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_1x1_ALGO_2) {

  1066.     require_compute_capability(7, 5);

  1067.     conv_bias::check_conv_bias(

  1068.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  1069.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.1f},

  1070.             handle_cuda(),

  1071.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma8x32x16",

  1072.             param::ConvBias::Format::CHWN4,

  1073.             conv_bias::get_int8_chwn4_args_small_batch(1));

  1074. }

  1075. 

  1076. 

  1077. #if CUDA_VERSION >= 10020

  1078. /// \note: we only check several cases and block sizes in megdnn_test, the full

  1079. /// testcases are written in cutlass repository

  1080. TEST_F(CUDA, CUTLASS_CONV_BIAS_INT8_NCHW32_IMMA) {

  1081.     require_compute_capability_eq(7, 5);

  1082.     Checker<ConvBiasForward> checker(handle_cuda());

  1083.     auto check = [&checker](const std::string& algo) {

  1084.         checker.set_before_exec_callback(

  1085.                 conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(algo.c_str()));

  1086.         UniformIntRNG rng{-8, 8};

  1087.         UniformIntRNG bias_rng{-50, 50};

  1088.         UniformIntRNG const_rng{1, 1};

  1089.         // use scale that are all integers to avoid rouding error

  1090.         checker.set_rng(0, &rng)

  1091.                 .set_rng(1, &rng)

  1092.                 .set_rng(2, &bias_rng)

  1093.                 .set_rng(3, &rng)

  1094.                 .set_dtype(0, dtype::QuantizedS8{6.0f})

  1095.                 .set_dtype(1, dtype::QuantizedS8{1.0f})

  1096.                 .set_dtype(2, dtype::QuantizedS32{6.0f})

  1097.                 .set_dtype(3, dtype::QuantizedS8{1.0f})

  1098.                 .set_dtype(4, dtype::QuantizedS8{6.0f})

  1099.                 .set_epsilon(1e-3);

  1100.         param::ConvBias param;

  1101.         param.pad_h = param.pad_w = 1;

  1102.         param.stride_h = param.stride_w = 1;

  1103.         param.format = param::ConvBias::Format::NCHW32;

  1104.         checker.set_param(param).execs({{16, 16, 7, 7, 32},

  1105.                                         {512, 16, 3, 3, 32},

  1106.                                         {1, 16, 1, 1, 32},

  1107.                                         {},

  1108.                                         {}});

  1109.         param.nonlineMode = param::ConvBias::NonlineMode::RELU;

  1110.         checker.set_param(param).execs({{16, 16, 7, 7, 32},

  1111.                                         {512, 16, 1, 1, 32},

  1112.                                         {1, 16, 1, 1, 32},

  1113.                                         {},

  1114.                                         {}});

  1115.         param.nonlineMode = param::ConvBias::NonlineMode::H_SWISH;

  1116.         checker.set_param(param).execs({{16, 16, 7, 7, 32},

  1117.                                         {512, 16, 3, 3, 32},

  1118.                                         {1, 16, 1, 1, 32},

  1119.                                         {},

  1120.                                         {}});

  1121.         // use non integer scale

  1122.         param.nonlineMode = param::ConvBias::NonlineMode::H_SWISH;

  1123.         checker.set_dtype(0, dtype::QuantizedS8{1.1f})

  1124.                 .set_dtype(1, dtype::QuantizedS8{1.2f})

  1125.                 .set_dtype(2, dtype::QuantizedS32{1.1f * 1.2f})

  1126.                 .set_dtype(3, dtype::QuantizedS8{1.1f})

  1127.                 .set_dtype(4, dtype::QuantizedS8{6.0f})

  1128.                 .set_epsilon(1 + 1e-3)

  1129.                 .set_max_avg_error(1e-1)

  1130.                 .set_max_avg_biased_error(1e-1)

  1131.                 .execs({{16, 16, 7, 7, 32},

  1132.                         {512, 16, 3, 3, 32},

  1133.                         {1, 16, 1, 1, 32},

  1134.                         {16, 16, 7, 7, 32},

  1135.                         {}});

  1136.     };

  1137.     std::string algo = ConvBias::algo_name<ConvBias::DirectParam>(

  1138.             "INT8_NCHW32_IMMA_IMPLICIT_GEMM_256X128X64_64X64X64",

  1139.             ConvBias::DirectParam{});

  1140.     check(algo);

  1141.     algo = ConvBias::algo_name<ConvBias::DirectParam>(

  1142.             "INT8_NCHW32_IMMA_IMPLICIT_GEMM_32X64X64_32X16X64",

  1143.             ConvBias::DirectParam{});

  1144.     check(algo);

  1145. }

  1146. #endif

  1147. 

  1148. #if MEGDNN_WITH_BENCHMARK

  1149. TEST_F(CUDA, BENCHMARK_CONV_BIAS_INT8_CHWN4) {

  1150.     require_compute_capability(6, 1);

  1151.     benchmark_target_algo(

  1152.             handle_cuda(), get_resnet50_bench_args(), dtype::QuantizedS8{1.2f},

  1153.             dtype::QuantizedS8{1.3f}, dtype::QuantizedS32{1.2f * 1.3f},

  1154.             dtype::QuantizedS8{1.0f}, "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  1155.             param::ConvBias::Format::CHWN4);

  1156. }

  1157. 

  1158. TEST_F(CUDA, BENCHMARK_CONV_BIAS_INT8_NCHW4) {

  1159.     require_compute_capability(6, 1);

  1160.     benchmark_target_algo(

  1161.             handle_cuda(), get_resnet50_bench_args(), dtype::QuantizedS8{1.2f},

  1162.             dtype::QuantizedS8{1.3f}, dtype::QuantizedS32{1.2f * 1.3f},

  1163.             dtype::QuantizedS8{1.0f}, "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

  1164.             param::ConvBias::Format::NCHW4);

  1165. }

  1166. 

  1167. TEST_F(CUDA, BENCHMARK_CONV_BIAS_INT8_CHWN4_TENSORCORE) {

  1168.     require_compute_capability(7, 5);

  1169.     benchmark_target_algo_with_cudnn_tsc(

  1170.             handle_cuda(), get_resnet50_bench_args(256),

  1171.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  1172.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.0f},

  1173.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  1174.             param::ConvBias::Format::CHWN4);

  1175. }

  1176. 

  1177. TEST_F(CUDA, BENCHMARK_CONV_BIAS_INT8_CHWN4_TENSORCORE_ALL_ALGO) {

  1178.     require_compute_capability(7, 5);

  1179.     benchmark_target_algo_with_cudnn_tsc(

  1180.             handle_cuda(), get_resnet50_bench_args(256),

  1181.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  1182.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.0f}, nullptr,

  1183.             param::ConvBias::Format::CHWN4);

  1184. }

  1185. 

  1186. TEST_F(CUDA, BENCHMARK_CONV_BIAS_INT8_CHWN4_DET_ALL_ALGO) {

  1187.     require_compute_capability(7, 5);

  1188.     benchmark_target_algo_with_cudnn_tsc(

  1189.             handle_cuda(), get_detection_bench_args(), dtype::QuantizedS8{1.2f},

  1190.             dtype::QuantizedS8{1.3f}, dtype::QuantizedS32{1.2f * 1.3f},

  1191.             dtype::QuantizedS8{1.0f}, nullptr, param::ConvBias::Format::CHWN4);

  1192. }

  1193. 

  1194. TEST_F(CUDA, BENCHMARK_CONV_BIAS_INT8_NCHW4_TENSORCORE) {

  1195.     require_compute_capability(7, 5);

  1196.     benchmark_target_algo_with_cudnn_tsc(

  1197.             handle_cuda(), get_resnet50_bench_args(256),

  1198.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  1199.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.0f},

  1200.             "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  1201.             param::ConvBias::Format::NCHW4);

  1202. }

  1203. 

  1204. TEST_F(CUDA, BENCHMARK_CONV_BIAS_INT8_CHWN4_SMALL_CHANNEL) {

  1205.     require_compute_capability(6, 1);

  1206.     std::vector<BenchArgs> args;

  1207.     args.push_back(BenchArgs{64, 4, 224, 224, 64, 7, 2});

  1208.     benchmark_target_algo(

  1209.             handle_cuda(), args, dtype::QuantizedS8{1.2f},

  1210.             dtype::QuantizedS8{1.3f}, dtype::QuantizedS32{1.2f * 1.3f},

  1211.             dtype::QuantizedS8{1.0f}, "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  1212.             param::ConvBias::Format::CHWN4);

  1213. }

  1214. 

  1215. 

  1216. #if CUDA_VERSION >= 10020

  1217. TEST_F(CUDA, BENCHMARK_CUTLASS_CONV_BIAS_INT8_NCHW32) {

  1218.     require_compute_capability(7, 5);

  1219.     benchmark_target_algo_with_cudnn_tsc(

  1220.             handle_cuda(), get_resnet50_bench_args(256),

  1221.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  1222.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.0f},

  1223.             "DIRECT:INT8_NCHW32_IMMA_IMPLICIT_GEMM",

  1224.             param::ConvBias::Format::NCHW32);

  1225. }

  1226. #endif

  1227. 

  1228. TEST_F(CUDA, BENCHMARK_CUTLASS_CONV_BIAS_INT8_NCHW4) {

  1229.     require_compute_capability(6, 1);

  1230.     benchmark_target_algo(

  1231.             handle_cuda(), get_resnet50_bench_args(64),

  1232.             dtype::QuantizedS8{1.2f}, dtype::QuantizedS8{1.3f},

  1233.             dtype::QuantizedS32{1.2f * 1.3f}, dtype::QuantizedS8{1.0f},

  1234.             "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM", param::ConvBias::Format::NCHW4);

  1235. }

  1236. #endif

  1237. }  // namespace test

  1238. }  // namespace megdnn

  1239. 

  1240. #undef V1

  1241. #undef V

  1242. 

  1243. // vim: syntax=cpp.doxygen