inference-engine/thirdparty/clDNN/tests/test_cases/custom_gpu_primitive_test.cpp

   1 /*
   2 // Copyright (c) 2016 Intel Corporation
   3 //
   4 // Licensed under the Apache License, Version 2.0 (the "License");
   5 // you may not use this file except in compliance with the License.
   6 // You may obtain a copy of the License at
   7 //
   8 //      http://www.apache.org/licenses/LICENSE-2.0
   9 //
  10 // Unless required by applicable law or agreed to in writing, software
  11 // distributed under the License is distributed on an "AS IS" BASIS,
  12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
  13 // See the License for the specific language governing permissions and
  14 // limitations under the License.
  15 */
  16
  17 ///////////////////////////////////////////////////////////////////////////////////////////////////
  18 #include <gtest/gtest.h>
  19 #include "api/CPP/memory.hpp"
  20 #include <api/CPP/input_layout.hpp>
  21 #include "api/CPP/eltwise.hpp"
  22 #include "api/CPP/reorder.hpp"
  23 #include "api/CPP/custom_gpu_primitive.hpp"
  24 #include <api/CPP/engine.hpp>
  25 #include <api/CPP/topology.hpp>
  26 #include <api/CPP/network.hpp>
  27 #include <api/CPP/engine.hpp>
  28 #include "test_utils/test_utils.h"
  29
  30
  31 namespace cldnn
  32 {
  33         template<> struct type_to_data_type<FLOAT16> { static const data_types value = data_types::f16; };
  34 }
  35
  36 using namespace cldnn;
  37 using namespace tests;
  38
  39 TEST(custom_gpu_primitive_f32, add_basic_in2x2x2x2) {
  40     //  Input2   : 2x2x2
  41     //  Input  : 2x2x2x2
  42     //  Output : 2x2x2x2
  43
  44     //  Input:
  45     //  f0: b0:  1    2  b1:   0    0
  46     //  f0: b0:  3    4  b1:   0.5 -0.5
  47     //  f1: b0:  5    6  b1:   1.5  5.2
  48     //  f1: b0:  7    8  b1:   12   8
  49     //
  50     //  Input2
  51     //  f0: b0: 0.5  5   b1: 2.5  7
  52     //  f0: b0: 15  -2   b1: 17   6.5
  53     //  f1: b0: 0.5  2   b1: 2.5  4
  54     //  f1: b0: 8   -0.5 b1: 10   -2.5
  55     //
  56     //  Output:
  57     //  f0: b0:   1.5  7    b1:  2.5   7
  58     //  f0: b0:   18   2    b1:  17.5  6
  59     //  f1: b0:   5.5  8    b1:   4    9.2
  60     //  f1: b0:   15  16.5  b1:  22    16.5
  61     //
  62
  63     const auto& engine = get_test_engine();
  64
  65     auto input = memory::allocate(engine, { data_types::f32, format::yxfb, { 2, 2, 2, 2 } });
  66     auto input2 = memory::allocate(engine, { data_types::f32, format::yxfb, { 2, 2, 2, 2 } });
  67
  68     std::string kernel_code =
  69         R"__krnl(
  70             __kernel void add_kernel(const __global float* input0, const __global float* input1, __global float* output)
  71             {
  72                 const unsigned idx = get_global_id(0);
  73                 output[idx] = input0[idx] + input1[idx];
  74             }
  75         )__krnl";
  76     std::string entry_point = "add_kernel";
  77     std::vector<cldnn_arg> parameters = { {arg_input, 0}, {arg_input, 1 }, {arg_output, 0 } };
  78     layout output_layout = { data_types::f32, format::yxfb,{ 2, 2, 2, 2 } };
  79     std::vector<size_t> gws = { output_layout.count() };
  80     topology topology;
  81     topology.add(input_layout("input", input.get_layout()));
  82     topology.add(input_layout("input2", input2.get_layout()));
  83     topology.add(custom_gpu_primitive(
  84         "user_kernel",
  85         { "input", "input2" },
  86         { kernel_code },
  87         entry_point,
  88         parameters,
  89         "-cl-mad-enable",
  90         output_layout,
  91         gws));
  92
  93     set_values(input, {
  94         1.f,   0.f, 5.f, 1.5f,
  95         2.f,   0.f, 6.f, 5.2f,
  96         3.f,  0.5f, 7.f, 12.f,
  97         4.f, -0.5f, 8.f,  8.f
  98     });
  99
 100     set_values(input2, {
 101         0.5f,   2.5f,  0.5f,  2.5f,
 102          5.f,   7.f,    2.f,   4.f,
 103         15.f,  17.f,    8.f,  10.f,
 104         -2.f,  6.5f,  -0.5f, -2.5f });
 105
 106     network network(engine, topology);
 107
 108     network.set_input_data("input", input);
 109     network.set_input_data("input2", input2);
 110     auto outputs = network.execute();
 111
 112     EXPECT_EQ(outputs.size(), size_t(1));
 113     EXPECT_EQ(outputs.begin()->first, "user_kernel");
 114
 115     auto output = outputs.at("user_kernel").get_memory();
 116
 117     float answers[16] = { 1.5f, 2.5f,   5.5f,    4.f,
 118                           7.f,   7.f,    8.f,   9.2f,
 119                           18.f,17.5f,   15.f,   22.f,
 120                           2.f,   6.f,   7.5f,  5.5f };
 121
 122     auto output_ptr = output.pointer<float>();
 123
 124     for (int i = 0; i < 16; i++)
 125     {
 126         EXPECT_TRUE(are_equal(answers[i], output_ptr[i]));
 127     }
 128 }
 129
 130 template<data_types DType>
 131 void add_basic_in2x2x2x2_with_reorder()
 132 {
 133     //  Input2   : 2x2x2
 134     //  Input  : 2x2x2x2
 135     //  Output : 2x2x2x2
 136
 137     //  Input:
 138     //  f0: b0:  1    2  b1:   0    0
 139     //  f0: b0:  3    4  b1:   0.5 -0.5
 140     //  f1: b0:  5    6  b1:   1.5  5.2
 141     //  f1: b0:  7    8  b1:   12   8
 142     //
 143     //  Input2
 144     //  f0: b0: 0.5  5   b1: 2.5  7
 145     //  f0: b0: 15  -2   b1: 17   6.5
 146     //  f1: b0: 0.5  2   b1: 2.5  4
 147     //  f1: b0: 8   -0.5 b1: 10   -2.5
 148     //
 149     //  Output:
 150     //  f0: b0:   1.5  7    b1:  2.5   7
 151     //  f0: b0:   18   2    b1:  17.5  6
 152     //  f1: b0:   5.5  8    b1:   4    9.2
 153     //  f1: b0:   15  16.5  b1:  22    16.5
 154     //
 155
 156     const auto& engine = get_test_engine();
 157
 158     auto input = memory::allocate(engine, { data_types::f32, format::yxfb,{ 2, 2, 2, 2 } });
 159     auto input2 = memory::allocate(engine, { data_types::f32, format::yxfb,{ 2, 2, 2, 2 } });
 160
 161     std::string data_type_string = "float";
 162     switch (DType)
 163     {
 164     case data_types::i32:
 165         data_type_string = "int";
 166         break;
 167     case data_types::i64:
 168         data_type_string = "long";
 169         break;
 170     default:
 171         throw std::runtime_error("Test does not support this data format!");
 172     }
 173
 174     std::string kernel_code =
 175         "__kernel void add_kernel(const __global " + data_type_string + "* input0, const __global " + data_type_string + "* input1, __global " + data_type_string + "* output)\n" +
 176         "   {\n" +
 177         "       const unsigned idx = get_global_id(0);\n" +
 178         "       output[idx] = input0[idx] + input1[idx];\n" +
 179         "   }\n";
 180     std::string entry_point = "add_kernel";
 181     std::vector<cldnn_arg> parameters = { { arg_input, 0 },{ arg_input, 1 },{ arg_output, 0 } };
 182     layout output_layout = { DType, format::yxfb,{ 2, 2, 2, 2 } };
 183     std::vector<size_t> gws = { output_layout.count() };
 184     topology topology;
 185     topology.add(input_layout("input", input.get_layout()));
 186     topology.add(input_layout("input2", input2.get_layout()));
 187     topology.add(reorder("to_int1", "input", { DType, format::yxfb,{ 2,2,2,2 } }));
 188     topology.add(reorder("to_int2", "input2", { DType, format::yxfb,{ 2,2,2,2 } }));
 189     topology.add(custom_gpu_primitive(
 190         "user_kernel",
 191         { "to_int1", "to_int2" },
 192         { kernel_code },
 193         entry_point,
 194         parameters,
 195         "-cl-mad-enable",
 196         output_layout,
 197         gws));
 198     topology.add(reorder("to_float", "user_kernel", { data_types::f32, format::yxfb,{ 2, 2, 2, 2 } }));
 199
 200     set_values(input, {
 201         1.f,   0.f, 5.f, 1.f,
 202         2.f,   0.f, 6.f, 5.f,
 203         3.f,  0.f, 7.f, 12.f,
 204         4.f,  0.f, 8.f,  8.f
 205     });
 206
 207     set_values(input2, {
 208         0.f,   2.f,  0.f,  2.f,
 209         5.f,   7.f,    2.f,   4.f,
 210         15.f,  17.f,    8.f,  10.f,
 211         -2.f,  6.f,  0.f, -2.f });
 212
 213     network network(engine, topology);
 214
 215     network.set_input_data("input", input);
 216     network.set_input_data("input2", input2);
 217     auto outputs = network.execute();
 218
 219     ASSERT_EQ(outputs.size(), size_t(1));
 220     EXPECT_EQ(outputs.begin()->first, "to_float");
 221
 222     auto output = outputs.at("to_float").get_memory();
 223
 224     float answers[16] = { 1.f, 2.f,   5.f,    3.f,
 225         7.f,   7.f,    8.f,   9.f,
 226         18.f,17.f,   15.f,   22.f,
 227         2.f,   6.f,   8.f,  6.f };
 228
 229     auto output_ptr = output.pointer<float>();
 230
 231     for (int i = 0; i < 16; i++)
 232     {
 233         EXPECT_TRUE(are_equal(answers[i], output_ptr[i]));
 234     }
 235 }
 236
 237 TEST(custom_gpu_primitive_int32, add_basic_in2x2x2x2) {
 238     add_basic_in2x2x2x2_with_reorder<data_types::i32>();
 239 }
 240
 241 TEST(custom_gpu_primitive_int64, add_basic_in2x2x2x2) {
 242     add_basic_in2x2x2x2_with_reorder<data_types::i64>();
 243 }
 244
 245 TEST(custom_gpu_primitive_f32, eltwise_add_basic_in2x2x2x2) {
 246     //  Input2   : 2x2x2
 247     //  Input  : 2x2x2x2
 248     //  Output : 2x2x2x2
 249
 250     //  Input:
 251     //  f0: b0:  1    2  b1:   0    0
 252     //  f0: b0:  3    4  b1:   0.5 -0.5
 253     //  f1: b0:  5    6  b1:   1.5  5.2
 254     //  f1: b0:  7    8  b1:   12   8
 255     //
 256     //  Input2
 257     //  f0: b0: 0.5  5   b1: 2.5  7
 258     //  f0: b0: 15  -2   b1: 17   6.5
 259     //  f1: b0: 0.5  2   b1: 2.5  4
 260     //  f1: b0: 8   -0.5 b1: 10   -2.5
 261     //
 262     //  Output:
 263     //  f0: b0:   1.5  7    b1:  2.5   7
 264     //  f0: b0:   18   2    b1:  17.5  6
 265     //  f1: b0:   5.5  8    b1:   4    9.2
 266     //  f1: b0:   15  16.5  b1:  22    16.5
 267     //
 268
 269     const auto& engine = get_test_engine();
 270
 271     auto input = memory::allocate(engine, { data_types::f32, format::yxfb,{ 2, 2, 2, 2 } });
 272     auto input2 = memory::allocate(engine, { data_types::f32, format::yxfb,{ 2, 2, 2, 2 } });
 273
 274     std::string kernel_code =
 275         R"__krnl(
 276             __kernel void add_kernel(const __global float* input0, __global float* output)
 277             {
 278                 const unsigned idx = get_global_id(0);
 279                 output[idx] = input0[idx] + 1;
 280             }
 281         )__krnl";
 282     std::string entry_point = "add_kernel";
 283     std::vector<cldnn_arg> parameters = { { arg_input, 0 },{ arg_output, 0 } };
 284     layout output_layout = { data_types::f32, format::yxfb,{ 2, 2, 2, 2 } };
 285     std::vector<size_t> gws = { output_layout.count() };
 286     topology topology;
 287     topology.add(input_layout("input", input.get_layout()));
 288     topology.add(input_layout("input2", input2.get_layout()));
 289     topology.add(eltwise("eltwise", {"input", "input2"}, eltwise_mode::sum));
 290     topology.add(custom_gpu_primitive(
 291         "user_kernel",
 292         { "eltwise" },
 293         { kernel_code },
 294         entry_point,
 295         parameters,
 296         "-cl-mad-enable",
 297         output_layout,
 298         gws));
 299
 300     set_values(input, {
 301         1.f,   0.f, 5.f, 1.5f,
 302         2.f,   0.f, 6.f, 5.2f,
 303         3.f,  0.5f, 7.f, 12.f,
 304         4.f, -0.5f, 8.f,  8.f
 305     });
 306
 307     set_values(input2, {
 308         0.5f,   2.5f,  0.5f,  2.5f,
 309         5.f,   7.f,    2.f,   4.f,
 310         15.f,  17.f,    8.f,  10.f,
 311         -2.f,  6.5f,  -0.5f, -2.5f });
 312
 313     network network(engine, topology);
 314
 315     network.set_input_data("input", input);
 316     network.set_input_data("input2", input2);
 317     auto outputs = network.execute();
 318
 319     EXPECT_EQ(outputs.size(), size_t(1));
 320     EXPECT_EQ(outputs.begin()->first, "user_kernel");
 321
 322     auto output = outputs.at("user_kernel").get_memory();
 323
 324     float answers[16] =
 325       { 2.5f,  3.5f,  6.5f,   5.f,
 326          8.f,   8.f,   9.f, 10.2f,
 327         19.f, 18.5f,  16.f,  23.f,
 328          3.f,   7.f,  8.5f,  6.5f };
 329
 330     auto output_ptr = output.pointer<float>();
 331
 332     for (int i = 0; i < 16; i++)
 333     {
 334         EXPECT_TRUE(are_equal(answers[i], output_ptr[i]));
 335     }
 336 }
 337
 338 TEST(custom_gpu_primitive_f32, add_eltwise_basic_in2x2x2x2) {
 339     //  Input2   : 2x2x2
 340     //  Input  : 2x2x2x2
 341     //  Output : 2x2x2x2
 342
 343     //  Input:
 344     //  f0: b0:  1    2  b1:   0    0
 345     //  f0: b0:  3    4  b1:   0.5 -0.5
 346     //  f1: b0:  5    6  b1:   1.5  5.2
 347     //  f1: b0:  7    8  b1:   12   8
 348     //
 349     //  Input2
 350     //  f0: b0: 0.5  5   b1: 2.5  7
 351     //  f0: b0: 15  -2   b1: 17   6.5
 352     //  f1: b0: 0.5  2   b1: 2.5  4
 353     //  f1: b0: 8   -0.5 b1: 10   -2.5
 354     //
 355     //  Output:
 356     //  f0: b0:   1.5  7    b1:  2.5   7
 357     //  f0: b0:   18   2    b1:  17.5  6
 358     //  f1: b0:   5.5  8    b1:   4    9.2
 359     //  f1: b0:   15  16.5  b1:  22    16.5
 360     //
 361
 362     const auto& engine = get_test_engine();
 363
 364     auto input = memory::allocate(engine, { data_types::f32, format::yxfb,{ 2, 2, 2, 2 } });
 365     auto input2 = memory::allocate(engine, { data_types::f32, format::yxfb,{ 2, 2, 2, 2 } });
 366
 367     std::string kernel_code =
 368         R"__krnl(
 369             __kernel void add_kernel(const __global float* input0, __global float* output)
 370             {
 371                 const unsigned idx = get_global_id(0);
 372                 output[idx] = input0[idx] + SCALAR;
 373             }
 374         )__krnl";
 375     std::string entry_point = "add_kernel";
 376     std::vector<cldnn_arg> parameters = { { arg_input, 0 },{ arg_output, 0 } };
 377     layout output_layout = { data_types::f32, format::yxfb,{ 2, 2, 2, 2 } };
 378     std::vector<size_t> gws = { output_layout.count() };
 379     topology topology;
 380     topology.add(input_layout("input", input.get_layout()));
 381     topology.add(input_layout("input2", input2.get_layout()));
 382     topology.add(custom_gpu_primitive(
 383         "user_kernel",
 384         { "input" },
 385         { kernel_code },
 386         entry_point,
 387         parameters,
 388         "-cl-mad-enable -DSCALAR=1",
 389         output_layout,
 390         gws));
 391     topology.add(eltwise("eltwise", {"user_kernel", "input2"}, eltwise_mode::sum));
 392
 393     set_values(input, {
 394         1.f,   0.f, 5.f, 1.5f,
 395         2.f,   0.f, 6.f, 5.2f,
 396         3.f,  0.5f, 7.f, 12.f,
 397         4.f, -0.5f, 8.f,  8.f
 398     });
 399
 400     set_values(input2, {
 401         0.5f,   2.5f,  0.5f,  2.5f,
 402         5.f,   7.f,    2.f,   4.f,
 403         15.f,  17.f,    8.f,  10.f,
 404         -2.f,  6.5f,  -0.5f, -2.5f });
 405
 406     network network(engine, topology);
 407
 408     network.set_input_data("input", input);
 409     network.set_input_data("input2", input2);
 410     auto outputs = network.execute();
 411
 412     EXPECT_EQ(outputs.size(), size_t(1));
 413     EXPECT_EQ(outputs.begin()->first, "eltwise");
 414
 415     auto output = outputs.at("eltwise").get_memory();
 416
 417     float answers[16] =
 418     { 2.5f,  3.5f,  6.5f,   5.f,
 419         8.f,   8.f,   9.f, 10.2f,
 420         19.f, 18.5f,  16.f,  23.f,
 421         3.f,   7.f,  8.5f,  6.5f };
 422
 423     auto output_ptr = output.pointer<float>();
 424
 425     for (int i = 0; i < 16; i++)
 426     {
 427         EXPECT_TRUE(are_equal(answers[i], output_ptr[i]));
 428     }
 429 }
 430
 431 TEST(custom_gpu_primitive_f32, two_kernels_with_same_entry_point_basic_in2x2x2x2) {
 432     //  Input2   : 2x2x2
 433     //  Input  : 2x2x2x2
 434     //  Output : 2x2x2x2
 435
 436     //  Input:
 437     //  f0: b0:  1    2  b1:   0    0
 438     //  f0: b0:  3    4  b1:   0.5 -0.5
 439     //  f1: b0:  5    6  b1:   1.5  5.2
 440     //  f1: b0:  7    8  b1:   12   8
 441     //
 442     //  Input2
 443     //  f0: b0: 0.5  5   b1: 2.5  7
 444     //  f0: b0: 15  -2   b1: 17   6.5
 445     //  f1: b0: 0.5  2   b1: 2.5  4
 446     //  f1: b0: 8   -0.5 b1: 10   -2.5
 447     //
 448     //  Output:
 449     //  f0: b0:   1.5  7    b1:  2.5   7
 450     //  f0: b0:   18   2    b1:  17.5  6
 451     //  f1: b0:   5.5  8    b1:   4    9.2
 452     //  f1: b0:   15  16.5  b1:  22    16.5
 453     //
 454
 455     const auto& engine = get_test_engine();
 456
 457     auto input = memory::allocate(engine, { data_types::f32, format::yxfb,{ 2, 2, 2, 2 } });
 458
 459     std::string kernel_code1 =
 460         R"__krnl(
 461             __kernel void add_kernel(const __global float* input0, __global float* output)
 462             {
 463                 const unsigned idx = get_global_id(0);
 464                 output[idx] = input0[idx] + SCALAR;
 465             }
 466         )__krnl";
 467
 468     std::string kernel_code2 =
 469         R"__krnl(
 470             __kernel void add_kernel(const __global float* input0, __global float* output)
 471             {
 472                 const unsigned idx = get_global_id(0);
 473                 output[idx] = input0[idx] + 2*SCALAR;
 474             }
 475         )__krnl";
 476     std::string entry_point = "add_kernel";
 477     std::vector<cldnn_arg> parameters = { { arg_input, 0 },{ arg_output, 0 } };
 478     layout output_layout = { data_types::f32, format::yxfb,{ 2, 2, 2, 2 } };
 479     std::vector<size_t> gws = { output_layout.count() };
 480     topology topology;
 481     topology.add(input_layout("input", input.get_layout()));
 482     topology.add(custom_gpu_primitive(
 483         "user_kernel1",
 484         { "input" },
 485         { kernel_code1 },
 486         entry_point,
 487         parameters,
 488         "-cl-mad-enable -DSCALAR=1",
 489         output_layout,
 490         gws));
 491     topology.add(custom_gpu_primitive(
 492         "user_kernel2",
 493         { "user_kernel1" },
 494         { kernel_code2 },
 495         entry_point,
 496         parameters,
 497         "-cl-mad-enable -DSCALAR=3",
 498         output_layout,
 499         gws));
 500
 501     set_values(input, {
 502         1.f,   0.f, 5.f, 1.5f,
 503         2.f,   0.f, 6.f, 5.2f,
 504         3.f,  0.5f, 7.f, 12.f,
 505         4.f, -0.5f, 8.f,  8.f
 506     });
 507
 508     network network(engine, topology);
 509
 510     network.set_input_data("input", input);
 511     auto outputs = network.execute();
 512
 513     EXPECT_EQ(outputs.size(), size_t(1));
 514     EXPECT_EQ(outputs.begin()->first, "user_kernel2");
 515
 516     auto output = outputs.at("user_kernel2").get_memory();
 517
 518     auto output_ptr = output.pointer<float>();
 519     auto input_ptr = input.pointer<float>();
 520
 521     for (int i = 0; i < 16; i++)
 522     {
 523         EXPECT_TRUE(are_equal(input_ptr[i] + 7, output_ptr[i]));
 524     }
 525 }
 526
 527 TEST(custom_gpu_primitive_u8, add_basic_in2x2x2x2) {
 528     const auto& engine = get_test_engine();
 529
 530     auto input = memory::allocate(engine, { data_types::u8, format::yxfb,{ 2, 2, 2, 2 } });
 531     auto input2 = memory::allocate(engine, { data_types::u8, format::yxfb,{ 2, 2, 2, 2 } });
 532
 533     std::string kernel_code =
 534         R"__krnl(
 535             __kernel void add_kernel(const __global uchar* input0, const __global uchar* input1, __global uchar* output)
 536             {
 537                 const unsigned idx = get_global_id(0);
 538                 output[idx] = input0[idx] + input1[idx];
 539             }
 540         )__krnl";
 541     std::string entry_point = "add_kernel";
 542     std::vector<cldnn_arg> parameters = { { arg_input, 0 },{ arg_input, 1 },{ arg_output, 0 } };
 543     layout output_layout = { data_types::u8, format::yxfb,{ 2, 2, 2, 2 } };
 544     std::vector<size_t> gws = { output_layout.count() };
 545     topology topology;
 546     topology.add(input_layout("input", input.get_layout()));
 547     topology.add(input_layout("input2", input2.get_layout()));
 548     topology.add(custom_gpu_primitive(
 549         "user_kernel",
 550         { "input", "input2" },
 551         { kernel_code },
 552         entry_point,
 553         parameters,
 554         "-cl-mad-enable",
 555         output_layout,
 556         gws));
 557
 558     set_values<unsigned char>(input, {
 559           1,   0,   5,    1,
 560         200, 100, 160,  150,
 561         130,   0, 175,   12,
 562           4, 100,   8,  180
 563     });
 564
 565     set_values<unsigned char>(input2, {
 566          0,  2,  0,  2,
 567         55, 75, 20,  4,
 568         15, 17, 80, 10,
 569          2, 60,  0, 20
 570     });
 571
 572     network network(engine, topology);
 573
 574     network.set_input_data("input", input);
 575     network.set_input_data("input2", input2);
 576     auto outputs = network.execute();
 577
 578     EXPECT_EQ(outputs.size(), size_t(1));
 579     EXPECT_EQ(outputs.begin()->first, "user_kernel");
 580
 581     auto output = outputs.at("user_kernel").get_memory();
 582
 583     unsigned char answers[16] = {
 584           1,   2,   5,   3,
 585         255, 175, 180, 154,
 586         145,  17, 255,  22,
 587           6, 160,   8, 200
 588     };
 589
 590     auto output_ptr = output.pointer<unsigned char>();
 591
 592     for (int i = 0; i < 16; i++)
 593     {
 594         EXPECT_TRUE(are_equal(answers[i], output_ptr[i]));
 595     }
 596 }