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: 8d248a6a9a
Branches Tags
${ item.name }
Create branch ${ searchTerm }
from '8d248a6a9a'
${ noResults }
MegEngine/dnn/test/cuda/conv_bias_int8.cpp

1114 lines
48 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-2021 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. #include "test/common/tensor.h"

  22. #include "test/common/workspace_wrapper.h"

  23. #include "test/cuda/conv_test_utils.h"

  24. 

  25. 

  26. 

  27. namespace megdnn {

  28. namespace test {

  29. namespace conv{

  30. 

  31. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_CUDNN_CONVOLUTION) {

  32.     require_compute_capability(7, 5);

  33.     conv_bias::check_conv_bias(

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

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

  36.             handle_cuda(), "DEFAULT:CUDNN:ConvBiasActivation:",

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

  38. }

  39. 

  40. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_1x1) {

  41.     require_compute_capability(6, 1);

  42.     conv_bias::check_conv_bias(

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

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

  45.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

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

  47. }

  48. 

  49. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_3x3) {

  50.     require_compute_capability(6, 1);

  51.     conv_bias::check_conv_bias(

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

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

  54.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

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

  56. }

  57. 

  58. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_5x5) {

  59.     require_compute_capability(6, 1);

  60.     conv_bias::check_conv_bias(

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

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

  63.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

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

  65. }

  66. 

  67. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_7x7) {

  68.     require_compute_capability(6, 1);

  69.     conv_bias::check_conv_bias(

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

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

  72.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

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

  74. }

  75. 

  76. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_WITH_Z) {

  77.     require_compute_capability(6, 1);

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

  79.     checker.set_before_exec_callback(

  80.             conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  81.                     "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM"));

  82.     UniformIntRNG rng{-3, 3};

  83.     UniformIntRNG bias_rng{-50, 50};

  84.     checker.set_rng(0, &rng)

  85.             .set_rng(1, &rng)

  86.             .set_rng(2, &bias_rng)

  87.             .set_rng(3, &rng)

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

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

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

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

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

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

  94.             .set_max_avg_error(1e-1)

  95.             .set_max_avg_biased_error(1e-1);

  96.     param::ConvBias param;

  97.     param.pad_h = param.pad_w = 1;

  98.     param.stride_h = param.stride_w = 1;

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

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

  101.                                     {16, 4, 3, 3, 4},

  102.                                     {1, 4, 1, 1, 4},

  103.                                     {32, 4, 12, 12, 4},

  104.                                     {}});

  105. }

  106. 

  107. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_STRIDE2_WITH_Z) {

  108.     require_compute_capability(6, 1);

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

  110.     checker.set_before_exec_callback(

  111.             conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  112.                     "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM"));

  113.     UniformIntRNG rng{-3, 3};

  114.     UniformIntRNG bias_rng{-50, 50};

  115.     checker.set_rng(0, &rng)

  116.             .set_rng(1, &rng)

  117.             .set_rng(2, &bias_rng)

  118.             .set_rng(3, &rng)

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

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

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

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

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

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

  125.             .set_max_avg_error(1e-1)

  126.             .set_max_avg_biased_error(1e-1);

  127.     param::ConvBias param;

  128.     param.pad_h = param.pad_w = 1;

  129.     param.stride_h = param.stride_w = 2;

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

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

  132.                                     {16, 4, 3, 3, 4},

  133.                                     {1, 4, 1, 1, 4},

  134.                                     {32, 4, 6, 6, 4},

  135.                                     {}});

  136. }

  137. 

  138. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_CHECK_BOUNDS_1x1) {

  139.     require_compute_capability(6, 1);

  140.     conv_bias::check_conv_bias(

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

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

  143.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

  144.             param::ConvBias::Format::NCHW4,

  145.             conv_bias::get_int8_nchw4_args_check_bounds(1));

  146. }

  147. 

  148. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_CHECK_BOUNDS_3x3) {

  149.     require_compute_capability(6, 1);

  150.     conv_bias::check_conv_bias(

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

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

  153.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

  154.             param::ConvBias::Format::NCHW4,

  155.             conv_bias::get_int8_nchw4_args_check_bounds(3));

  156. }

  157. 

  158. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_CHECK_BOUNDS_5x5) {

  159.     require_compute_capability(6, 1);

  160.     conv_bias::check_conv_bias(

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

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

  163.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

  164.             param::ConvBias::Format::NCHW4,

  165.             conv_bias::get_int8_nchw4_args_check_bounds(5));

  166. }

  167. 

  168. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_CHECK_BOUNDS_7x7) {

  169.     require_compute_capability(6, 1);

  170.     conv_bias::check_conv_bias(

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

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

  173.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

  174.             param::ConvBias::Format::NCHW4,

  175.             conv_bias::get_int8_nchw4_args_check_bounds(7));

  176. }

  177. 

  178. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4) {

  179.     require_compute_capability(6, 1);

  180.     conv_bias::check_conv_bias(

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

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

  183.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

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

  185. }

  186. 

  187. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_WITH_Z) {

  188.     require_compute_capability(6, 1);

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

  190.     checker.set_before_exec_callback(

  191.             conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  192.                     "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM"));

  193.     UniformIntRNG rng{-3, 3};

  194.     UniformIntRNG bias_rng{-50, 50};

  195.     checker.set_rng(0, &rng)

  196.             .set_rng(1, &rng)

  197.             .set_rng(2, &bias_rng)

  198.             .set_rng(3, &rng)

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

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

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

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

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

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

  205.             .set_max_avg_error(1e-1)

  206.             .set_max_avg_biased_error(1e-1);

  207.     param::ConvBias param;

  208.     param.pad_h = param.pad_w = 1;

  209.     param.stride_h = param.stride_w = 1;

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

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

  212.                                     {4, 3, 3, 16, 4},

  213.                                     {4, 1, 1, 1, 4},

  214.                                     {4, 12, 12, 32, 4},

  215.                                     {}});

  216. }

  217. 

  218. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_HSWISH) {

  219.     require_compute_capability(6, 1);

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

  221.     checker.set_before_exec_callback(

  222.             conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  223.                     "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM"));

  224.     UniformIntRNG rng{-3, 3};

  225.     UniformIntRNG bias_rng{-50, 50};

  226.     checker.set_rng(0, &rng)

  227.             .set_rng(1, &rng)

  228.             .set_rng(2, &bias_rng)

  229.             .set_rng(3, &rng)

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

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

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

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

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

  235.             .set_max_avg_error(1e-1)

  236.             .set_max_avg_biased_error(1e-1);

  237.     param::ConvBias param;

  238.     param.pad_h = param.pad_w = 1;

  239.     param.stride_h = param.stride_w = 1;

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

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

  242.     checker.set_param(param).execs(

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

  244. }

  245. 

  246. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_CHECK_BOUNDS) {

  247.     require_compute_capability(6, 1);

  248.     conv_bias::check_conv_bias(

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

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

  251.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  252.             param::ConvBias::Format::CHWN4,

  253.             conv_bias::get_int8_chwn4_args_check_bounds(3));

  254. }

  255. 

  256. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_SMALL_CHANNEL_1x1) {

  257.     require_compute_capability(6, 1);

  258.     conv_bias::check_conv_bias(

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

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

  261.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  262.             param::ConvBias::Format::CHWN4,

  263.             conv_bias::get_int8_chwn4_small_channel_args(1));

  264. }

  265. 

  266. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_SMALL_CHANNEL_3x3) {

  267.     require_compute_capability(6, 1);

  268.     conv_bias::check_conv_bias(

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

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

  271.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  272.             param::ConvBias::Format::CHWN4,

  273.             conv_bias::get_int8_chwn4_small_channel_args(3));

  274. }

  275. 

  276. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_SMALL_CHANNEL_5x5) {

  277.     require_compute_capability(6, 1);

  278.     conv_bias::check_conv_bias(

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

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

  281.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  282.             param::ConvBias::Format::CHWN4,

  283.             conv_bias::get_int8_chwn4_small_channel_args(5));

  284. }

  285. 

  286. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_SMALL_CHANNEL_7x7) {

  287.     require_compute_capability(6, 1);

  288.     conv_bias::check_conv_bias(

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

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

  291.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  292.             param::ConvBias::Format::CHWN4,

  293.             conv_bias::get_int8_chwn4_small_channel_args(7));

  294. }

  295. 

  296. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_SMALL_CHANNEL_CHECK_BOUNDS) {

  297.     require_compute_capability(6, 1);

  298.     conv_bias::check_conv_bias(

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

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

  301.             handle_cuda(), "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM",

  302.             param::ConvBias::Format::NCHW4,

  303.             conv_bias::get_int8_nchw4_small_channel_args_check_bounds(3));

  304. }

  305. 

  306. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_SMALL_CHANNEL_1x1_CHECK_BOUNDS) {

  307.     require_compute_capability(6, 1);

  308.     conv_bias::check_conv_bias(

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

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

  311.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  312.             param::ConvBias::Format::CHWN4,

  313.             conv_bias::get_int8_chwn4_small_channel_args_check_bounds(1));

  314. }

  315. 

  316. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_SMALL_CHANNEL_5x5_CHECK_BOUNDS) {

  317.     require_compute_capability(6, 1);

  318.     conv_bias::check_conv_bias(

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

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

  321.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  322.             param::ConvBias::Format::CHWN4,

  323.             conv_bias::get_int8_chwn4_small_channel_args_check_bounds(5));

  324. }

  325. 

  326. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_SMALL_CHANNEL_7x7_CHECK_BOUNDS) {

  327.     require_compute_capability(6, 1);

  328.     conv_bias::check_conv_bias(

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

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

  331.             handle_cuda(), "INT8_CHWN4_DOTPROD_IMPLICIT_GEMM",

  332.             param::ConvBias::Format::CHWN4,

  333.             conv_bias::get_int8_chwn4_small_channel_args_check_bounds(7));

  334. }

  335. 

  336. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_TENSORCORE_1x1) {

  337.     require_compute_capability(7, 5);

  338.     conv_bias::check_conv_bias(

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

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

  341.             handle_cuda(), "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  342.             param::ConvBias::Format::NCHW4,

  343.             conv_bias::get_int8_nchw4_tensorcore_args(1));

  344. }

  345. 

  346. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_TENSORCORE_3x3) {

  347.     require_compute_capability(7, 5);

  348.     conv_bias::check_conv_bias(

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

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

  351.             handle_cuda(), "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  352.             param::ConvBias::Format::NCHW4,

  353.             conv_bias::get_int8_nchw4_tensorcore_args(3));

  354. }

  355. 

  356. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_TENSORCORE_5x5) {

  357.     require_compute_capability(7, 5);

  358.     conv_bias::check_conv_bias(

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

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

  361.             handle_cuda(), "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  362.             param::ConvBias::Format::NCHW4,

  363.             conv_bias::get_int8_nchw4_tensorcore_args(5));

  364. }

  365. 

  366. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_TENSORCORE_7x7) {

  367.     require_compute_capability(7, 5);

  368.     conv_bias::check_conv_bias(

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

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

  371.             handle_cuda(), "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  372.             param::ConvBias::Format::NCHW4,

  373.             conv_bias::get_int8_nchw4_tensorcore_args(7));

  374. }

  375. 

  376. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_TENSORCORE_CHECK_BOUNDS_ALGO_0) {

  377.     require_compute_capability(7, 5);

  378.     conv_bias::check_conv_bias(

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

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

  381.             handle_cuda(), "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  382.             param::ConvBias::Format::NCHW4,

  383.             conv_bias::get_int8_nchw4_args_check_bounds(3));

  384. }

  385. 

  386. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_TENSORCORE_CHECK_BOUNDS_ALGO_1) {

  387.     require_compute_capability(7, 5);

  388.     conv_bias::check_conv_bias(

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

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

  391.             handle_cuda(), "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma8x32x16",

  392.             param::ConvBias::Format::NCHW4,

  393.             conv_bias::get_int8_nchw4_args_check_bounds(3));

  394. }

  395. 

  396. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_TENSORCORE_CHECK_BOUNDS_ALGO_2) {

  397.     require_compute_capability(7, 5);

  398.     conv_bias::check_conv_bias(

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

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

  401.             handle_cuda(), "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma32x8x16",

  402.             param::ConvBias::Format::NCHW4,

  403.             conv_bias::get_int8_nchw4_args_check_bounds(3));

  404. }

  405. 

  406. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_TENSORCORE_ALGO_0) {

  407.     require_compute_capability(7, 5);

  408.     conv_bias::check_conv_bias(

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

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

  411.             handle_cuda(), "INT8_CHWN4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  412.             param::ConvBias::Format::CHWN4,

  413.             conv_bias::get_int8_chwn4_tensorcore_args(3));

  414. }

  415. 

  416. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_TENSORCORE_ALGO_1) {

  417.     require_compute_capability(7, 5);

  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.1f},

  421.             handle_cuda(), "INT8_CHWN4_IMMA_IMPLICIT_GEMM_mma32x8x16",

  422.             param::ConvBias::Format::CHWN4,

  423.             conv_bias::get_int8_chwn4_tensorcore_args(3));

  424. }

  425. 

  426. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_TENSORCORE_ALGO_2) {

  427.     require_compute_capability(7, 5);

  428.     conv_bias::check_conv_bias(

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

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

  431.             handle_cuda(), "INT8_CHWN4_IMMA_IMPLICIT_GEMM_mma8x32x16",

  432.             param::ConvBias::Format::CHWN4,

  433.             conv_bias::get_int8_chwn4_tensorcore_args(3));

  434. }

  435. 

  436. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_TENSORCORE_CHECK_BOUNDS_1x1) {

  437.     require_compute_capability(7, 5);

  438.     conv_bias::check_conv_bias(

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

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

  441.             handle_cuda(), "INT8_CHWN4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  442.             param::ConvBias::Format::CHWN4,

  443.             conv_bias::get_int8_chwn4_args_check_bounds(1));

  444. }

  445. 

  446. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_TENSORCORE_CHECK_BOUNDS_5x5) {

  447.     require_compute_capability(7, 5);

  448.     conv_bias::check_conv_bias(

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

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

  451.             handle_cuda(), "INT8_CHWN4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  452.             param::ConvBias::Format::CHWN4,

  453.             conv_bias::get_int8_chwn4_args_check_bounds(5));

  454. }

  455. 

  456. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_TENSORCORE_CHECK_BOUNDS_7x7) {

  457.     require_compute_capability(7, 5);

  458.     conv_bias::check_conv_bias(

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

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

  461.             handle_cuda(), "INT8_CHWN4_IMMA_IMPLICIT_GEMM_mma16x16x16",

  462.             param::ConvBias::Format::CHWN4,

  463.             conv_bias::get_int8_chwn4_args_check_bounds(7));

  464. }

  465. 

  466. TEST_F(CUDA, CONV_BIAS_INT8_NCHW4_TENSORCORE_WITH_Z) {

  467.     require_compute_capability(7, 5);

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

  469.     checker.set_before_exec_callback(

  470.             conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  471.                     "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma16x16x16"));

  472.     UniformIntRNG rng{-3, 3};

  473.     UniformIntRNG bias_rng{-50, 50};

  474.     checker.set_rng(0, &rng)

  475.             .set_rng(1, &rng)

  476.             .set_rng(2, &bias_rng)

  477.             .set_rng(3, &rng)

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

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

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

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

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

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

  484.             .set_max_avg_error(1e-1)

  485.             .set_max_avg_biased_error(1e-1);

  486.     param::ConvBias param;

  487.     param.pad_h = param.pad_w = 1;

  488.     param.stride_h = param.stride_w = 1;

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

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

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

  492.                                     {1, 16, 1, 1, 4},

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

  494.                                     {}});

  495. }

  496. 

  497. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_TENSORCORE_WITH_Z) {

  498.     require_compute_capability(7, 5);

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

  500.     checker.set_before_exec_callback(

  501.             conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  502.                     "INT8_CHWN4_IMMA_IMPLICIT_GEMM_mma16x16x16"));

  503.     UniformIntRNG rng{-3, 3};

  504.     UniformIntRNG bias_rng{-50, 50};

  505.     checker.set_rng(0, &rng)

  506.             .set_rng(1, &rng)

  507.             .set_rng(2, &bias_rng)

  508.             .set_rng(3, &rng)

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

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

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

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

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

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

  515.             .set_max_avg_error(1e-1)

  516.             .set_max_avg_biased_error(1e-1);

  517.     param::ConvBias param;

  518.     param.pad_h = param.pad_w = 1;

  519.     param.stride_h = param.stride_w = 1;

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

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

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

  523.                                     {16, 1, 1, 1, 4},

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

  525.                                     {}});

  526. }

  527. 

  528. TEST_F(CUDA,

  529.        CONV_BIAS_INT8_CHWN4_REFORMAT_FILTER_TENSORCORE_CHECK_BOUNDS_ALGO_0) {

  530.     require_compute_capability(7, 5);

  531.     conv_bias::check_conv_bias(

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

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

  534.             handle_cuda(),

  535.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_REORDER_FILTER_mma16x16x16",

  536.             param::ConvBias::Format::CHWN4,

  537.             conv_bias::get_int8_chwn4_args_check_bounds(3));

  538. }

  539. 

  540. TEST_F(CUDA,

  541.        CONV_BIAS_INT8_CHWN4_REFORMAT_FILTER_TENSORCORE_CHECK_BOUNDS_ALGO_1) {

  542.     require_compute_capability(7, 5);

  543.     conv_bias::check_conv_bias(

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

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

  546.             handle_cuda(),

  547.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_REORDER_FILTER_mma8x32x16",

  548.             param::ConvBias::Format::CHWN4,

  549.             conv_bias::get_int8_chwn4_args_check_bounds(3));

  550. }

  551. 

  552. TEST_F(CUDA,

  553.        CONV_BIAS_INT8_CHWN4_REFORMAT_FILTER_TENSORCORE_CHECK_BOUNDS_ALGO_2) {

  554.     require_compute_capability(7, 5);

  555.     conv_bias::check_conv_bias(

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

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

  558.             handle_cuda(),

  559.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_REORDER_FILTER_mma32x8x16",

  560.             param::ConvBias::Format::CHWN4,

  561.             conv_bias::get_int8_chwn4_args_check_bounds(3));

  562. }

  563. 

  564. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_REFORMAT_FILTER_TENSORCORE_ALGO_0) {

  565.     require_compute_capability(7, 5);

  566.     conv_bias::check_conv_bias(

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

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

  569.             handle_cuda(),

  570.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_REORDER_FILTER_mma16x16x16",

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

  572. }

  573. 

  574. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_REFORMAT_FILTER_TENSORCORE_ALGO_1) {

  575.     require_compute_capability(7, 5);

  576.     conv_bias::check_conv_bias(

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

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

  579.             handle_cuda(),

  580.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_REORDER_FILTER_mma8x32x16",

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

  582. }

  583. 

  584. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_REFORMAT_FILTER_TENSORCORE_ALGO_2) {

  585.     require_compute_capability(7, 5);

  586.     conv_bias::check_conv_bias(

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

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

  589.             handle_cuda(),

  590.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_REORDER_FILTER_mma32x8x16",

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

  592. }

  593. 

  594. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_ALGO_0) {

  595.     require_compute_capability(7, 5);

  596.     conv_bias::check_conv_bias(

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

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

  599.             handle_cuda(),

  600.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma16x16x16",

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

  602. }

  603. 

  604. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_ALGO_1) {

  605.     require_compute_capability(7, 5);

  606.     conv_bias::check_conv_bias(

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

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

  609.             handle_cuda(),

  610.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma8x32x16",

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

  612. }

  613. 

  614. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_ALGO_2) {

  615.     require_compute_capability(7, 5);

  616.     conv_bias::check_conv_bias(

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

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

  619.             handle_cuda(),

  620.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma32x8x16",

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

  622. }

  623. 

  624. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_1x1) {

  625.     require_compute_capability(7, 5);

  626.     conv_bias::check_conv_bias(

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

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

  629.             handle_cuda(),

  630.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma16x16x16",

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

  632. }

  633. 

  634. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_5x5) {

  635.     require_compute_capability(7, 5);

  636.     conv_bias::check_conv_bias(

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

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

  639.             handle_cuda(),

  640.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma16x16x16",

  641.             param::ConvBias::Format::CHWN4,

  642.             conv_bias::get_int8_chwn4_args_small_batch(5));

  643. }

  644. 

  645. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_7x7) {

  646.     require_compute_capability(7, 5);

  647.     conv_bias::check_conv_bias(

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

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

  650.             handle_cuda(),

  651.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma16x16x16",

  652.             param::ConvBias::Format::CHWN4,

  653.             conv_bias::get_int8_chwn4_args_small_batch(7));

  654. }

  655. 

  656. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_5x5_ALGO_1) {

  657.     require_compute_capability(7, 5);

  658.     conv_bias::check_conv_bias(

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

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

  661.             handle_cuda(),

  662.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma32x8x16",

  663.             param::ConvBias::Format::CHWN4,

  664.             conv_bias::get_int8_chwn4_args_small_batch(5));

  665. }

  666. 

  667. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_5x5_ALGO_2) {

  668.     require_compute_capability(7, 5);

  669.     conv_bias::check_conv_bias(

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

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

  672.             handle_cuda(),

  673.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma8x32x16",

  674.             param::ConvBias::Format::CHWN4,

  675.             conv_bias::get_int8_chwn4_args_small_batch(5));

  676. }

  677. 

  678. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_1x1_ALGO_1) {

  679.     require_compute_capability(7, 5);

  680.     conv_bias::check_conv_bias(

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

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

  683.             handle_cuda(),

  684.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma32x8x16",

  685.             param::ConvBias::Format::CHWN4,

  686.             conv_bias::get_int8_chwn4_args_small_batch(1));

  687. }

  688. 

  689. TEST_F(CUDA, CONV_BIAS_INT8_CHWN4_UNROLL_WIDTH_TENSORCORE_1x1_ALGO_2) {

  690.     require_compute_capability(7, 5);

  691.     conv_bias::check_conv_bias(

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

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

  694.             handle_cuda(),

  695.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_UNROLL_WIDTH_mma8x32x16",

  696.             param::ConvBias::Format::CHWN4,

  697.             conv_bias::get_int8_chwn4_args_small_batch(1));

  698. }

  699. 

  700. 

  701. TEST_F(CUDA, FALLBACK_CONV_QS8) {

  702.     require_compute_capability_eq(7, 5);

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

  704.     auto check = [&checker](const std::string&& algo,

  705.                             const std::string&& sub_algo) {

  706.         checker.set_before_exec_callback(

  707.                 conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  708.                         {algo.c_str(), {sub_algo.c_str()}}));

  709.         UniformIntRNG rng{-3, 3};

  710.         UniformIntRNG bias_rng{-50, 50};

  711.         checker.set_rng(0, &rng)

  712.                 .set_rng(1, &rng)

  713.                 .set_rng(2, &bias_rng)

  714.                 .set_rng(3, &rng)

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

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

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

  718.                 .set_dtype(3, dtype::QuantizedS8{19.990229f})

  719.                 .set_dtype(4, dtype::QuantizedS8{19.990228f})

  720.                 .set_epsilon(1e-3)

  721.                 .set_max_avg_error(1e-1)

  722.                 .set_max_avg_biased_error(1e-3);

  723.         param::ConvBias param;

  724.         param.pad_h = param.pad_w = 1;

  725.         param.stride_h = param.stride_w = 2;

  726.         param.format = param::ConvBias::Format::NCHW;

  727.         checker.set_param(param).execs({{16, 15, 14, 14},

  728.                                         {28, 15, 3, 3},

  729.                                         {1, 28, 1, 1},

  730.                                         {16, 28, 7, 7},

  731.                                         {}});

  732.         checker.set_param(param).execs({{16, 32, 14, 14},

  733.                                         {32, 32, 3, 3},

  734.                                         {1, 32, 1, 1},

  735.                                         {},

  736.                                         {}});

  737.     };

  738.     check("FALLBACK_CONV_NCHW_QS8", "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM");

  739. }

  740. 

  741. TEST_F(CUDA, FALLBACK_CONV_QS8_F32) {

  742.     require_compute_capability_eq(7, 5);

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

  744.     auto check = [&checker](const std::string&& algo,

  745.                             const std::string&& sub_algo) {

  746.         checker.set_before_exec_callback(

  747.                 conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  748.                         {algo.c_str(), {sub_algo.c_str()}}));

  749.         UniformIntRNG rng{-3, 3};

  750.         UniformFloatRNG bias_rng{-50.f, 50.f};

  751.         checker.set_rng(0, &rng)

  752.                 .set_rng(1, &rng)

  753.                 .set_rng(2, &bias_rng)

  754.                 .set_rng(3, &rng)

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

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

  757.                 .set_dtype(2, dtype::Float32{})

  758.                 .set_dtype(3, dtype::Float32{})

  759.                 .set_dtype(4, dtype::Float32{})

  760.                 .set_epsilon(1e-3)

  761.                 .set_max_avg_error(1e-1)

  762.                 .set_max_avg_biased_error(1e-3);

  763.         param::ConvBias param;

  764.         param.pad_h = param.pad_w = 1;

  765.         param.stride_h = param.stride_w = 2;

  766.         param.format = param::ConvBias::Format::NCHW;

  767.         checker.set_param(param).execs({{16, 15, 14, 14},

  768.                                         {28, 15, 3, 3},

  769.                                         {1, 28, 1, 1},

  770.                                         {16, 28, 7, 7},

  771.                                         {}});

  772.         checker.set_param(param).execs({{16, 32, 14, 14},

  773.                                         {32, 32, 3, 3},

  774.                                         {1, 32, 1, 1},

  775.                                         {},

  776.                                         {}});

  777.     };

  778.     check("FALLBACK_CONV_NCHW_QS8", "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM");

  779. }

  780. 

  781. TEST_F(CUDA, CUTLASS_CONV_BIAS_INT8_WEIGHT_PREPROCESS) {

  782.     require_compute_capability(6, 1);

  783.     Checker<ConvBiasForward, OprWeightPreprocessProxy<ConvBiasForward>> checker(

  784.             handle_cuda());

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

  786.         checker.set_before_exec_callback(

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

  788.         UniformIntRNG rng{-16, 16};

  789.         UniformIntRNG bias_rng{-50, 50};

  790.         UniformIntRNG const_rng{1, 1};

  791.         checker.set_rng(0, &rng)

  792.                 .set_rng(1, &rng)

  793.                 .set_rng(2, &bias_rng)

  794.                 .set_rng(3, &rng)

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

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

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

  798.                 .set_dtype(3, dtype::QuantizedS8{1.3f})

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

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

  801.                 .set_max_avg_error(1e-1)

  802.                 .set_max_avg_biased_error(1e-3);

  803.         param::ConvBias param;

  804.         param.pad_h = param.pad_w = 1;

  805.         param.stride_h = param.stride_w = 2;

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

  807.         checker.set_param(param).execs({{16, 4, 14, 14, 4},

  808.                                         {16, 4, 3, 3, 4},

  809.                                         {1, 4, 1, 1, 4},

  810.                                         {},

  811.                                         {}});

  812.     };

  813.     check("INT8_NCHW4_DOTPROD_IMPLICIT_GEMM_128X32X32_64X32X32");

  814.     check("INT8_NCHW4_DOTPROD_IMPLICIT_GEMM_16X64X8_16X64X8");

  815. }

  816. 

  817. #if CUDA_VERSION >= 10020

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

  819. /// full testcases are written in cutlass repository

  820. TEST_F(CUDA, CUTLASS_CONV_BIAS_INT8_NCHW32_IMMA) {

  821.     require_compute_capability_eq(7, 5);

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

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

  824.         checker.set_before_exec_callback(

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

  826.         UniformIntRNG rng{-8, 8};

  827.         UniformIntRNG bias_rng{-50, 50};

  828.         UniformIntRNG const_rng{1, 1};

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

  830.         checker.set_rng(0, &rng)

  831.                 .set_rng(1, &rng)

  832.                 .set_rng(2, &bias_rng)

  833.                 .set_rng(3, &rng)

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

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

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

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

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

  839.                 .set_epsilon(1e-3);

  840.         param::ConvBias param;

  841.         param.pad_h = param.pad_w = 1;

  842.         param.stride_h = param.stride_w = 1;

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

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

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

  846.                                         {1, 16, 1, 1, 32},

  847.                                         {},

  848.                                         {}});

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

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

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

  852.                                         {1, 16, 1, 1, 32},

  853.                                         {},

  854.                                         {}});

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

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

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

  858.                                         {1, 16, 1, 1, 32},

  859.                                         {},

  860.                                         {}});

  861.         // use non integer scale

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

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

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

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

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

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

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

  869.                 .set_max_avg_error(1e-1)

  870.                 .set_max_avg_biased_error(1e-1)

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

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

  873.                         {1, 16, 1, 1, 32},

  874.                         {16, 16, 7, 7, 32},

  875.                         {}});

  876.     };

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

  878.             "INT8_NCHW32_IMMA_IMPLICIT_GEMM_256X128X64_64X64X64",

  879.             ConvBias::DirectParam{});

  880.     check(algo);

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

  882.             "INT8_NCHW32_IMMA_IMPLICIT_GEMM_32X64X64_32X16X64",

  883.             ConvBias::DirectParam{});

  884.     check(algo);

  885. }

  886. #endif

  887. 

  888. TEST_F(CUDA, CUTLASS_CONV_BIAS_INT8_NCHW4_NCHW) {

  889.     require_compute_capability(6, 1);

  890.     using namespace conv_bias;

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

  892.     UniformIntRNG int_rng{-3, 3};

  893.     UniformFloatRNG float_rng{-50, 50};

  894.     ConvBias::Param param;

  895.     param.format = ConvBias::Param::Format::NCHW4_NCHW;

  896.     param.nonlineMode = ConvBias::Param::NonlineMode::IDENTITY;

  897.     checker.set_before_exec_callback(

  898.             conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  899.                     "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM"));

  900.     checker.set_dtype(0, dtype::QuantizedS8(1.9980618f))

  901.             .set_dtype(1, dtype::QuantizedS8(1.9980927f))

  902.             .set_dtype(2, dtype::Float32())

  903.             .set_dtype(3, dtype::Float32())

  904.             .set_dtype(4, dtype::Float32())

  905.             .set_rng(0, &int_rng)

  906.             .set_rng(1, &int_rng)

  907.             .set_rng(2, &float_rng)

  908.             .set_rng(3, &float_rng)

  909.             .set_param(param);

  910. 

  911.     auto opr = handle_cuda()->create_operator<ConvBias>();

  912. 

  913.     auto run = [&](const TensorShapeArray& shapes) {

  914.         opr->param() = param;

  915.         TensorLayout dst_layout;

  916.         opr->deduce_layout({shapes[0], dtype::Float32()},

  917.                            {shapes[1], dtype::Float32()}, {}, {}, dst_layout);

  918.         checker.execs({shapes[0], shapes[1], shapes[2], dst_layout, {}});

  919.     };

  920. 

  921.     run({{16, 4, 23, 40, 4}, {20, 4, 3, 3, 4}, {1, 20, 1, 1}});

  922.     run({{16, 4, 92, 160, 4}, {24, 4, 3, 3, 4}, {1, 24, 1, 1}});

  923.     run({{16, 4, 92, 160, 4}, {20, 4, 3, 3, 4}, {1, 20, 1, 1}});

  924.     run({{16, 4, 92, 160, 4}, {16, 4, 3, 3, 4}, {1, 16, 1, 1}});

  925.     run({{16, 4, 92, 160, 4}, {8, 4, 3, 3, 4}, {1, 8, 1, 1}});

  926.     run({{16, 4, 46, 80, 4}, {4, 4, 3, 3, 4}, {1, 4, 1, 1}});

  927. }

  928. 

  929. TEST_F(CUDA, CUTLASS_CONV_BIAS_INT8_NCHW4_NCHW32) {

  930.     require_compute_capability(6, 1);

  931.     using namespace conv_bias;

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

  933.     UniformIntRNG int_rng{-3, 3};

  934.     UniformIntRNG bias_rng{-50, 50};

  935.     ConvBias::Param param;

  936.     param.format = ConvBias::Param::Format::NCHW4_NCHW32;

  937.     param.nonlineMode = ConvBias::Param::NonlineMode::IDENTITY;

  938.     checker.set_before_exec_callback(

  939.             conv_bias::ConvBiasAlgoChecker<ConvBiasForward>(

  940.                     "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM"));

  941.     checker.set_dtype(0, dtype::QuantizedS8(1.9980618f))

  942.             .set_dtype(1, dtype::QuantizedS8(1.9980927f))

  943.             .set_dtype(2, dtype::QuantizedS32(1.9980618f * 1.9980927f))

  944.             .set_dtype(3, dtype::QuantizedS8(1.9980618f))

  945.             .set_dtype(4, dtype::QuantizedS8(1.9980618f))

  946.             .set_rng(0, &int_rng)

  947.             .set_rng(1, &int_rng)

  948.             .set_rng(2, &bias_rng)

  949.             .set_rng(3, &int_rng)

  950.             .set_param(param);

  951.     auto run = [&](const TensorShapeArray& shapes) {

  952.         checker.execs({shapes[0], shapes[1], shapes[2], {}, {}});

  953.     };

  954. 

  955.     run({{16, 4, 23, 40, 4}, {32, 4, 3, 3, 4}, {1, 1, 1, 1, 32}});

  956.     run({{16, 4, 92, 160, 4}, {32, 4, 3, 3, 4}, {1, 1, 1, 1, 32}});

  957.     run({{16, 4, 46, 80, 4}, {32, 4, 3, 3, 4}, {1, 1, 1, 1, 32}});

  958. }

  959. 

  960. #if CUDA_VERSION >= 10020

  961. TEST_F(CUDA, CUTLASS_CONV_BIAS_INT8_NCHW32_NCHW4) {

  962.     require_compute_capability(7, 5);

  963.     using namespace conv_bias;

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

  965.     UniformIntRNG int_rng{-3, 3};

  966.     UniformIntRNG bias_rng{-50, 50};

  967.     ConvBias::Param param;

  968.     param.format = ConvBias::Param::Format::NCHW32_NCHW4;

  969.     param.nonlineMode = ConvBias::Param::NonlineMode::IDENTITY;

  970.     checker.set_before_exec_callback(conv_bias::ConvBiasAlgoChecker<

  971.                                      ConvBiasForward>(

  972.             ConvBias::algo_name<ConvBias::DirectParam>(

  973.                     "INT8_NCHW32_IMMA_IMPLICIT_GEMM_256X128X64_64X64X64",

  974.                     ConvBias::DirectParam{})

  975.                     .c_str()));

  976.     checker.set_dtype(0, dtype::QuantizedS8(1.9980618f))

  977.             .set_dtype(1, dtype::QuantizedS8(1.9980927f))

  978.             .set_dtype(2, dtype::QuantizedS32(1.9980618f * 1.9980927f))

  979.             .set_dtype(3, dtype::QuantizedS8(1.9980618f))

  980.             .set_dtype(4, dtype::QuantizedS8(1.9980618f))

  981.             .set_rng(0, &int_rng)

  982.             .set_rng(1, &int_rng)

  983.             .set_rng(2, &bias_rng)

  984.             .set_rng(3, &int_rng)

  985.             .set_param(param);

  986.     auto run = [&](const TensorShapeArray& shapes) {

  987.         checker.execs({shapes[0], shapes[1], shapes[2], {}, {}});

  988.     };

  989. 

  990.     run({{16, 2, 23, 40, 32}, {20, 2, 3, 3, 32}, {1, 5, 1, 1, 4}});

  991.     run({{16, 1, 92, 160, 32}, {24, 1, 3, 3, 32}, {1, 6, 1, 1, 4}});

  992.     run({{16, 2, 46, 80, 32}, {4, 2, 3, 3, 32}, {1, 1, 1, 1, 4}});

  993. }

  994. #endif

  995. 

  996. #if MEGDNN_WITH_BENCHMARK

  997. TEST_F(CUDA, BENCHMARK_CONV_BIAS_INT8_CHWN4) {

  998.     require_compute_capability(6, 1);

  999.     benchmark_target_algo(

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

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

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

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

  1004. }

  1005. 

  1006. TEST_F(CUDA, BENCHMARK_CONV_BIAS_INT8_NCHW4) {

  1007.     require_compute_capability(6, 1);

  1008.     benchmark_target_algo(

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

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

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

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

  1013. }

  1014. 

  1015. TEST_F(CUDA, BENCHMARK_CONV_BIAS_INT8_CHWN4_TENSORCORE) {

  1016.     require_compute_capability(7, 5);

  1017.     benchmark_target_algo_with_cudnn_tsc(

  1018.             handle_cuda(), get_resnet50_bench_args(256),

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

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

  1021.             "INT8_CHWN4_IMMA_IMPLICIT_GEMM_mma16x16x16",

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

  1023. }

  1024. 

  1025. TEST_F(CUDA, BENCHMARK_CONV_BIAS_INT8_CHWN4_TENSORCORE_ALL_ALGO) {

  1026.     require_compute_capability(7, 5);

  1027.     benchmark_target_algo_with_cudnn_tsc(

  1028.             handle_cuda(), get_resnet50_bench_args(256),

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

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

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

  1032. }

  1033. 

  1034. TEST_F(CUDA, BENCHMARK_CONV_BIAS_INT8_CHWN4_DET_ALL_ALGO) {

  1035.     require_compute_capability(7, 5);

  1036.     benchmark_target_algo_with_cudnn_tsc(

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

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

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

  1040. }

  1041. 

  1042. TEST_F(CUDA, BENCHMARK_CONV_BIAS_INT8_NCHW4_TENSORCORE) {

  1043.     require_compute_capability(7, 5);

  1044.     benchmark_target_algo_with_cudnn_tsc(

  1045.             handle_cuda(), get_resnet50_bench_args(256),

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

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

  1048.             "INT8_NCHW4_IMMA_IMPLICIT_GEMM_mma16x16x16",

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

  1050. }

  1051. 

  1052. TEST_F(CUDA, BENCHMARK_CONV_BIAS_INT8_CHWN4_SMALL_CHANNEL) {

  1053.     require_compute_capability(6, 1);

  1054.     std::vector<BenchArgs> args;

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

  1056.     benchmark_target_algo(

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

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

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

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

  1061. }

  1062. 

  1063. 

  1064. #if CUDA_VERSION >= 10020

  1065. TEST_F(CUDA, BENCHMARK_CUTLASS_CONV_BIAS_INT8_NCHW32) {

  1066.     require_compute_capability(7, 5);

  1067.     benchmark_target_algo_with_cudnn_tsc(

  1068.             handle_cuda(), get_resnet50_bench_args(256),

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

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

  1071.             "DIRECT:INT8_NCHW32_IMMA_IMPLICIT_GEMM",

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

  1073. }

  1074. #endif

  1075. 

  1076. TEST_F(CUDA, BENCHMARK_CUTLASS_CONV_BIAS_INT8_NCHW4) {

  1077.     require_compute_capability(6, 1);

  1078.     benchmark_target_algo(

  1079.             handle_cuda(), get_resnet50_bench_args(64),

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

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

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

  1083. }

  1084. 

  1085. TEST_F(CUDA, BENCHMARK_SASS_CONV_BIAS_INT8_NCHW4_DET_FIRST) {

  1086.     require_compute_capability(6, 1);

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

  1088.             "SASS_INT8_NCHW4_DOTPROD_IMPLICIT_GEMM_128X32_64",

  1089.             ConvBias::DirectParam{});

  1090.     benchmark_target_algo(handle_cuda(), get_det_first_bench_args(16),

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

  1092.                           dtype::QuantizedS32{1.2f * 1.3f},

  1093.                           dtype::QuantizedS8{1.0f}, algo.c_str(),

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

  1095. }

  1096. 

  1097. TEST_F(CUDA, BENCHMARK_CUTLASS_CONV_BIAS_INT8_NCHW4_DET_FIRST) {

  1098.     require_compute_capability(6, 1);

  1099.     benchmark_target_algo(

  1100.             handle_cuda(), get_det_first_bench_args(16),

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

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

  1103.             "INT8_NCHW4_DOTPROD_IMPLICIT_GEMM_16", param::ConvBias::Format::NCHW4);

  1104. }

  1105. 

  1106. #endif

  1107. }

  1108. }  // namespace test

  1109. }  // namespace megdnn

  1110. 

  1111. 

  1112. 

  1113. // vim: syntax=cpp.doxygen