path: root/test_conformance/device_execution/enqueue_block.cpp

//
// Copyright (c) 2017 The Khronos Group Inc.
// 
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
//    http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
//
#include <stdio.h>
#include <string.h>
#include "harness/testHarness.h"
#include "harness/typeWrappers.h"

#include <vector>

#include "procs.h"
#include "utils.h"
#include <time.h>


#ifdef CL_VERSION_2_0
extern int gWimpyMode;

// clang-format off
static const char* enqueue_simple_block[] = { R"(
    void block_fn(size_t tid, int mul, __global int* res)
    {
      res[tid] = mul * 7 - 21;
    }

    kernel void enqueue_simple_block(__global int* res)
    {
      int multiplier = 3;
      size_t tid = get_global_id(0);

      void (^kernelBlock)(void) = ^{ block_fn(tid, multiplier, res); };

      res[tid] = -1;
      queue_t def_q = get_default_queue();
      ndrange_t ndrange = ndrange_1D(1);
      int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);
      if (enq_res != CLK_SUCCESS) { res[tid] = -1; return; }
    }
)" };

static const char* enqueue_block_with_local_arg1[] = { R"(
    #define LOCAL_MEM_SIZE 10

    void block_fn_local_arg1(size_t tid, int mul, __global int* res, __local int* tmp)
    {
      for (int i = 0; i < LOCAL_MEM_SIZE; i++)
      {
        tmp[i] = mul * 7 - 21;
        res[tid] += tmp[i];
      }
      res[tid] += 2;
    }

    kernel void enqueue_block_with_local_arg1(__global int* res)
    {
      int multiplier = 3;
      size_t tid = get_global_id(0);

      void (^kernelBlock)(__local void*) = ^(__local void* buf){ block_fn_local_arg1(tid, multiplier, res, (local int*)buf); };

      res[tid] = -2;
      queue_t def_q = get_default_queue();
      ndrange_t ndrange = ndrange_1D(1);
      int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock, (uint)(LOCAL_MEM_SIZE*sizeof(int)));
      if (enq_res != CLK_SUCCESS) { res[tid] = -1; return; }
    }
)" };

static const char* enqueue_block_with_local_arg2[] = { R"(
    #define LOCAL_MEM_SIZE 10

    void block_fn_local_arg1(size_t tid, int mul, __global int* res, __local int* tmp1, __local float4* tmp2)
    {
      for (int i = 0; i < LOCAL_MEM_SIZE; i++)
      {
        tmp1[i]   = mul * 7 - 21;
        tmp2[i].x = (float)(mul * 7 - 21);
        tmp2[i].y = (float)(mul * 7 - 21);
        tmp2[i].z = (float)(mul * 7 - 21);
        tmp2[i].w = (float)(mul * 7 - 21);

        res[tid] += tmp1[i];
        res[tid] += (int)(tmp2[i].x+tmp2[i].y+tmp2[i].z+tmp2[i].w);
      }
      res[tid] += 2;
    }

    kernel void enqueue_block_with_local_arg2(__global int* res)
    {
      int multiplier = 3;
      size_t tid = get_global_id(0);

      void (^kernelBlock)(__local void*, __local void*) = ^(__local void* buf1, __local void* buf2)
        { block_fn_local_arg1(tid, multiplier, res, (local int*)buf1, (local float4*)buf2); };

      res[tid] = -2;
      queue_t def_q = get_default_queue();
      ndrange_t ndrange = ndrange_1D(1);
      int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock, (uint)(LOCAL_MEM_SIZE*sizeof(int)), (uint)(LOCAL_MEM_SIZE*sizeof(float4)));
      if (enq_res != CLK_SUCCESS) { res[tid] = -1; return; }
    }
)" };

static const char* enqueue_block_with_wait_list[] = { R"(
    #define BLOCK_SUBMITTED 1
    #define BLOCK_COMPLETED 2
    #define CHECK_SUCCESS   0

    kernel void enqueue_block_with_wait_list(__global int* res)
    {
      size_t tid = get_global_id(0);

      clk_event_t user_evt = create_user_event();

      res[tid] = BLOCK_SUBMITTED;
      queue_t def_q = get_default_queue();
      ndrange_t ndrange = ndrange_1D(1);
      clk_event_t block_evt;
      int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &user_evt, &block_evt,
      ^{
          res[tid] = BLOCK_COMPLETED;
       });
      if (enq_res != CLK_SUCCESS) { res[tid] = -1; return; }

      retain_event(block_evt);
      release_event(block_evt);

      //check block is not started
      if (res[tid] == BLOCK_SUBMITTED)
      {
        clk_event_t my_evt;
        enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &block_evt, &my_evt,
        ^{
           //check block is completed
           if (res[tid] == BLOCK_COMPLETED) res[tid] = CHECK_SUCCESS;
         });
        release_event(my_evt);
      }

      set_user_event_status(user_evt, CL_COMPLETE);

      release_event(user_evt);
      release_event(block_evt);
    }
)" };

static const char* enqueue_block_with_wait_list_and_local_arg[] = { R"(
    #define LOCAL_MEM_SIZE 10
    #define BLOCK_COMPLETED 1
    #define BLOCK_SUBMITTED 2
    #define BLOCK_STARTED   3
    #define CHECK_SUCCESS   0

    void block_fn_local_arg(size_t tid, int mul, __global int* res, __local int* tmp)
    {
      res[tid] = BLOCK_STARTED;
      for (int i = 0; i < LOCAL_MEM_SIZE; i++)
      {
        tmp[i] = mul * 7 - 21;
        res[tid] += tmp[i];
      }
      if (res[tid] == BLOCK_STARTED) res[tid] = BLOCK_COMPLETED;
    }

    kernel void enqueue_block_with_wait_list_and_local_arg(__global int* res)
    {
      int multiplier = 3;
      size_t tid = get_global_id(0);
      clk_event_t user_evt = create_user_event();

      res[tid] = BLOCK_SUBMITTED;
      queue_t def_q = get_default_queue();
      ndrange_t ndrange = ndrange_1D(1);
      clk_event_t block_evt;
      int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &user_evt, &block_evt,
        ^(__local void* buf) {
           block_fn_local_arg(tid, multiplier, res, (__local int*)buf);
         }, LOCAL_MEM_SIZE*sizeof(int));
      if (enq_res != CLK_SUCCESS) { res[tid] = -1; return; }

      retain_event(block_evt);
      release_event(block_evt);

      //check block is not started
      if (res[tid] == BLOCK_SUBMITTED)
      {
        clk_event_t my_evt;
        enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &block_evt, &my_evt,
        ^{
           //check block is completed
           if (res[tid] == BLOCK_COMPLETED) res[tid] = CHECK_SUCCESS;
         });
        release_event(my_evt);
      }

      set_user_event_status(user_evt, CL_COMPLETE);

      release_event(user_evt);
      release_event(block_evt);
    }
)" };

static const char* enqueue_block_get_kernel_work_group_size[] = { R"(
    void block_fn(size_t tid, int mul, __global int* res)
    {
      res[tid] = mul * 7 - 21;
    }

    kernel void enqueue_block_get_kernel_work_group_size(__global int* res)
    {
        int multiplier = 3;
        size_t tid = get_global_id(0);

        void (^kernelBlock)(void) = ^{ block_fn(tid, multiplier, res); };

        size_t local_work_size = get_kernel_work_group_size(kernelBlock);
        if (local_work_size <= 0){ res[tid] = -1; return; }
        size_t global_work_size = local_work_size * 4;

        res[tid] = -1;
        queue_t q1 = get_default_queue();
        ndrange_t ndrange = ndrange_1D(global_work_size, local_work_size);

        int enq_res = enqueue_kernel(q1, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);
        if (enq_res != CLK_SUCCESS) { res[tid] = -1; return; }
    }
)" };

static const char* enqueue_block_get_kernel_preferred_work_group_size_multiple[] = { R"(
    void block_fn(size_t tid, int mul, __global int* res)
    {
      res[tid] = mul * 7 - 21;
    }

    kernel void enqueue_block_get_kernel_preferred_work_group_size_multiple(__global int* res)
    {
        int multiplier = 3;
        size_t tid = get_global_id(0);

        void (^kernelBlock)(void) = ^{ block_fn(tid, multiplier, res); };

        size_t local_work_size = get_kernel_preferred_work_group_size_multiple(kernelBlock);
        if (local_work_size <= 0){ res[tid] = -1; return; }
        size_t global_work_size = local_work_size * 4;

        res[tid] = -1;
        queue_t q1 = get_default_queue();
        ndrange_t ndrange = ndrange_1D(global_work_size, local_work_size);

        int enq_res = enqueue_kernel(q1, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);
        if (enq_res != CLK_SUCCESS) { res[tid] = -1; return; }
    }
)" };

static const char* enqueue_block_capture_event_profiling_info_after_execution[] = {
    "#define MAX_GWS " STRINGIFY_VALUE(MAX_GWS) "\n"
    , R"(
    __global ulong value[MAX_GWS*2] = {0};

    void block_fn(size_t tid, __global int* res)
    {
        res[tid] = -2;
    }

    void check_res(size_t tid, const clk_event_t evt, __global int* res)
    {
        capture_event_profiling_info (evt, CLK_PROFILING_COMMAND_EXEC_TIME, &value[tid*2]);

        if (value[tid*2] > 0 && value[tid*2+1] > 0) res[tid] =  0;
        else                                        res[tid] = -4;
        release_event(evt);
    }

    kernel void enqueue_block_capture_event_profiling_info_after_execution(__global int* res)
    {
        size_t tid = get_global_id(0);

        res[tid] = -1;
        queue_t def_q = get_default_queue();
        ndrange_t ndrange = ndrange_1D(1);
        clk_event_t block_evt1;

        void (^kernelBlock)(void)  = ^{ block_fn (tid, res);                   };

        int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 0, NULL, &block_evt1, kernelBlock);
        if (enq_res != CLK_SUCCESS) { res[tid] = -1; return; }

        void (^checkBlock) (void)  = ^{ check_res(tid, block_evt1, res);      };

        enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &block_evt1, NULL, checkBlock);
        if (enq_res != CLK_SUCCESS) { res[tid] = -3; return; }
    }
)" };

static const char* enqueue_block_capture_event_profiling_info_before_execution[] = {
    "#define MAX_GWS " STRINGIFY_VALUE(MAX_GWS) "\n"
    , R"(
    __global ulong value[MAX_GWS*2] = {0};

    void block_fn(size_t tid, __global int* res)
    {
        res[tid] = -2;
    }

    void check_res(size_t tid, const ulong *value, __global int* res)
    {
        if (value[tid*2] > 0 && value[tid*2+1] > 0) res[tid] =  0;
        else                                        res[tid] = -4;
    }

    kernel void enqueue_block_capture_event_profiling_info_before_execution(__global int* res)
    {
        int multiplier = 3;
        size_t tid = get_global_id(0);
        clk_event_t user_evt = create_user_event();

        res[tid] = -1;
        queue_t def_q = get_default_queue();
        ndrange_t ndrange = ndrange_1D(1);
        clk_event_t block_evt1;
        clk_event_t block_evt2;

        void (^kernelBlock)(void)  = ^{ block_fn (tid, res);                   };

        int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &user_evt, &block_evt1, kernelBlock);
        if (enq_res != CLK_SUCCESS) { res[tid] = -1; return; }

        capture_event_profiling_info (block_evt1, CLK_PROFILING_COMMAND_EXEC_TIME, &value[tid*2]);

        set_user_event_status(user_evt, CL_COMPLETE);

        void (^checkBlock) (void)  = ^{ check_res(tid, &value, res);      };

        enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &block_evt1, &block_evt2, checkBlock);
        if (enq_res != CLK_SUCCESS) { res[tid] = -3; return; }

        release_event(user_evt);
        release_event(block_evt1);
        release_event(block_evt2);
    }
)" };

static const char* enqueue_block_with_barrier[] = { R"(
    void block_fn(size_t tid, int mul, __global int* res)
    {
      if (mul > 0) barrier(CLK_GLOBAL_MEM_FENCE);
      res[tid] = mul * 7 -21;
    }

    void loop_fn(size_t tid, int n, __global int* res)
    {
      while (n > 0)
      {
        barrier(CLK_GLOBAL_MEM_FENCE);
        res[tid] = 0;
        --n;
      }
    }

    kernel void enqueue_block_with_barrier(__global int* res)
    {
      int multiplier = 3;
      size_t tid = get_global_id(0);
      queue_t def_q = get_default_queue();
      res[tid] = -1;
      size_t n = 256;

      void (^kernelBlock)(void) = ^{ block_fn(tid, multiplier, res); };

      ndrange_t ndrange = ndrange_1D(n);
      int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, kernelBlock);
      if (enq_res != CLK_SUCCESS) { res[tid] = -1; return; }

      void (^loopBlock)(void) = ^{ loop_fn(tid, n, res); };

      enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, loopBlock);
      if (enq_res != CLK_SUCCESS) { res[tid] = -1; return; }
    }
)" };

static const char* enqueue_marker_with_block_event[] = { R"(
    #define BLOCK_COMPLETED 1
    #define BLOCK_SUBMITTED 2
    #define CHECK_SUCCESS   0

    kernel void enqueue_marker_with_block_event(__global int* res)
    {
      size_t tid = get_global_id(0);

      clk_event_t user_evt = create_user_event();

      res[tid] = BLOCK_SUBMITTED;
      queue_t def_q = get_default_queue();
      ndrange_t ndrange = ndrange_1D(1);

      clk_event_t block_evt1;
      clk_event_t marker_evt;

      int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &user_evt, &block_evt1,
      ^{
         res[tid] = BLOCK_COMPLETED;
       });
      if (enq_res != CLK_SUCCESS) { res[tid] = -2; return; }

      enq_res = enqueue_marker(def_q, 1, &block_evt1, &marker_evt);
      if (enq_res != CLK_SUCCESS) { res[tid] = -3; return; }

      retain_event(marker_evt);
      release_event(marker_evt);

      //check block is not started
      if (res[tid] == BLOCK_SUBMITTED)
      {
        clk_event_t my_evt;
        enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &marker_evt, &my_evt,
        ^{
           //check block is completed
           if (res[tid] == BLOCK_COMPLETED) res[tid] = CHECK_SUCCESS;
         });
        release_event(my_evt);
      }

      set_user_event_status(user_evt, CL_COMPLETE);

      release_event(block_evt1);
      release_event(marker_evt);
      release_event(user_evt);
    }
)" };

static const char* enqueue_marker_with_user_event[] = { R"(
    #define BLOCK_COMPLETED 1
    #define BLOCK_SUBMITTED 2
    #define CHECK_SUCCESS   0

    kernel void enqueue_marker_with_user_event(__global int* res)
    {
      size_t tid = get_global_id(0);
      uint multiplier = 7;

      clk_event_t user_evt = create_user_event();

      res[tid] = BLOCK_SUBMITTED;
      queue_t def_q = get_default_queue();
      ndrange_t ndrange = ndrange_1D(1);

      clk_event_t marker_evt;
      clk_event_t block_evt;

      int enq_res = enqueue_marker(def_q, 1, &user_evt, &marker_evt);
      if (enq_res != CLK_SUCCESS) { res[tid] = -1; return; }

      retain_event(marker_evt);
      release_event(marker_evt);

      enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &marker_evt, &block_evt,
      ^{
         if (res[tid] == BLOCK_SUBMITTED) res[tid] = CHECK_SUCCESS;
       });

      //check block is not started
      if (res[tid] != BLOCK_SUBMITTED)  { res[tid] = -2; return; }

      set_user_event_status(user_evt, CL_COMPLETE);

      release_event(block_evt);
      release_event(marker_evt);
      release_event(user_evt);
    }
)" };

static const char* enqueue_marker_with_mixed_events[] = { R"(
    #define BLOCK_COMPLETED 1
    #define BLOCK_SUBMITTED 2
    #define CHECK_SUCCESS   0

    kernel void enqueue_marker_with_mixed_events(__global int* res)
    {
      size_t tid = get_global_id(0);

      clk_event_t mix_ev[2];
      mix_ev[0] = create_user_event();

      res[tid] = BLOCK_SUBMITTED;
      queue_t def_q = get_default_queue();
      ndrange_t ndrange = ndrange_1D(1);

      int enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &mix_ev[0], &mix_ev[1],
      ^{
         res[tid] = BLOCK_COMPLETED;
       });
      if (enq_res != CLK_SUCCESS) { res[tid] = -2; return; }

      clk_event_t marker_evt;

      enq_res = enqueue_marker(def_q, 2, mix_ev, &marker_evt);
      if (enq_res != CLK_SUCCESS) { res[tid] = -3; return; }

      retain_event(marker_evt);
      release_event(marker_evt);

      //check block is not started
      if (res[tid] == BLOCK_SUBMITTED)
      {
        clk_event_t my_evt;
        enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &marker_evt, &my_evt,
        ^{
           //check block is completed
           if (res[tid] == BLOCK_COMPLETED) res[tid] = CHECK_SUCCESS;
         });
        release_event(my_evt);
      }

      set_user_event_status(mix_ev[0], CL_COMPLETE);

      release_event(mix_ev[1]);
      release_event(marker_evt);
      release_event(mix_ev[0]);
    }
)" };

static const char* enqueue_block_with_mixed_events[] = { R"(
    kernel void enqueue_block_with_mixed_events(__global int* res)
    {
      int enq_res;
      size_t tid = get_global_id(0);
      clk_event_t mix_ev[3];
      mix_ev[0] = create_user_event();
      queue_t def_q = get_default_queue();
      ndrange_t ndrange = ndrange_1D(1);
      res[tid] = -2;

      enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, 1, &mix_ev[0], &mix_ev[1], ^{ res[tid]++; });
      if (enq_res != CLK_SUCCESS) { res[tid] = -1; return; }

      enq_res = enqueue_marker(def_q, 1, &mix_ev[1], &mix_ev[2]);
      if (enq_res != CLK_SUCCESS) { res[tid] = -3; return; }

      enq_res = enqueue_kernel(def_q, CLK_ENQUEUE_FLAGS_NO_WAIT, ndrange, sizeof(mix_ev)/sizeof(mix_ev[0]), mix_ev, NULL, ^{ res[tid]++; });
      if (enq_res != CLK_SUCCESS) { res[tid] = -4; return; }

      set_user_event_status(mix_ev[0], CL_COMPLETE);

      release_event(mix_ev[0]);
      release_event(mix_ev[1]);
      release_event(mix_ev[2]);
    }
)" };
// clang-format on

static const kernel_src sources_enqueue_block[] =
{
    KERNEL(enqueue_simple_block),
    // Block with local mem
    KERNEL(enqueue_block_with_local_arg1),
    KERNEL(enqueue_block_with_local_arg2),
    KERNEL(enqueue_block_with_wait_list),
    KERNEL(enqueue_block_with_wait_list_and_local_arg),
    // WG size built-ins
    KERNEL(enqueue_block_get_kernel_work_group_size),
    KERNEL(enqueue_block_get_kernel_preferred_work_group_size_multiple),
    // Event profiling info
    KERNEL(enqueue_block_capture_event_profiling_info_after_execution),
    KERNEL(enqueue_block_capture_event_profiling_info_before_execution),
    // Marker
    KERNEL(enqueue_marker_with_block_event),
    KERNEL(enqueue_marker_with_user_event),
    // Mixed events
    KERNEL(enqueue_marker_with_mixed_events),
    KERNEL(enqueue_block_with_mixed_events),
    // Barrier
    KERNEL(enqueue_block_with_barrier),

};
static const size_t num_kernels_enqueue_block = arr_size(sources_enqueue_block);

static int check_kernel_results(cl_int* results, cl_int len)
{
    for(cl_int i = 0; i < len; ++i)
    {
        if(results[i] != 0) return i;
    }
    return -1;
}

int test_enqueue_block(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
    cl_uint i;
    cl_int n, err_ret, res = 0;
    clCommandQueueWrapper dev_queue;
    cl_int kernel_results[MAX_GWS] = {0};

    size_t ret_len;
    cl_uint max_queues = 1;
    cl_uint maxQueueSize = 0;
    err_ret = clGetDeviceInfo(device, CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE, sizeof(maxQueueSize), &maxQueueSize, 0);
    test_error(err_ret, "clGetDeviceInfo(CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE) failed");

    err_ret = clGetDeviceInfo(device, CL_DEVICE_MAX_ON_DEVICE_QUEUES, sizeof(max_queues), &max_queues, &ret_len);
    test_error(err_ret, "clGetDeviceInfo(CL_DEVICE_MAX_ON_DEVICE_QUEUES) failed");

    size_t max_local_size = 1;
    err_ret = clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(max_local_size), &max_local_size, &ret_len);
    test_error(err_ret, "clGetDeviceInfo(CL_DEVICE_MAX_WORK_GROUP_SIZE) failed");

    cl_queue_properties queue_prop_def[] =
    {
        CL_QUEUE_PROPERTIES, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE|CL_QUEUE_ON_DEVICE|CL_QUEUE_ON_DEVICE_DEFAULT|CL_QUEUE_PROFILING_ENABLE,
        CL_QUEUE_SIZE, maxQueueSize,
        0
    };

    dev_queue = clCreateCommandQueueWithProperties(context, device, queue_prop_def, &err_ret);
    test_error(err_ret, "clCreateCommandQueueWithProperties(CL_QUEUE_DEVICE|CL_QUEUE_DEFAULT) failed");

    size_t global_size = MAX_GWS;
    size_t local_size = (max_local_size > global_size/16) ? global_size/16 : max_local_size;
    if(gWimpyMode)
    {
        global_size = 4;
        local_size = 2;
    }

    size_t failCnt = 0;
    for(i = 0; i < num_kernels_enqueue_block; ++i)
    {
        if (!gKernelName.empty() && gKernelName != sources_enqueue_block[i].kernel_name)
            continue;

        log_info("Running '%s' kernel (%d of %d) ...\n", sources_enqueue_block[i].kernel_name, i + 1, num_kernels_enqueue_block);
        err_ret = run_n_kernel_args(context, queue, sources_enqueue_block[i].lines, sources_enqueue_block[i].num_lines, sources_enqueue_block[i].kernel_name, local_size, global_size, kernel_results, sizeof(kernel_results), 0, NULL);
        if(check_error(err_ret, "'%s' kernel execution failed", sources_enqueue_block[i].kernel_name)) { ++failCnt; res = -1; }
        else if((n = check_kernel_results(kernel_results, arr_size(kernel_results))) >= 0 && check_error(-1, "'%s' kernel results validation failed: [%d] returned %d expected 0", sources_enqueue_block[i].kernel_name, n, kernel_results[n])) res = -1;
        else log_info("'%s' kernel is OK.\n", sources_enqueue_block[i].kernel_name);
    }

    if (failCnt > 0)
    {
      log_error("ERROR: %d of %d kernels failed.\n", failCnt, num_kernels_enqueue_block);
    }

    return res;
}



#endif