OpenCL-CTS/test_conformance/device_execution/execute_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 "../../test_common/harness/testHarness.h"
#include "../../test_common/harness/typeWrappers.h"

#include <vector>

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


#ifdef CL_VERSION_2_0

static const char* block_global_scope[] =
{
    NL, "int __constant globalVar = 7;"
    NL, "int (^__constant globalBlock)(int) = ^int(int num)"
    NL, "{"
    NL, "   return globalVar * num * (1+ get_global_id(0));"
    NL, "};"
    NL, "kernel void block_global_scope(__global int* res)"
    NL, "{"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  res[tid] = globalBlock(3) - 21*(tid + 1);"
    NL, "}"
    NL
};

static const char* block_kernel_scope[] =
{
    NL, "kernel void block_kernel_scope(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 3;"
    NL, "  int (^kernelBlock)(int) = ^(int num)"
    NL, "  {"
    NL, "    return num * multiplier;"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  multiplier = 8;"
    NL, "  res[tid] = kernelBlock(7) - 21;"
    NL, "}"
    NL
};

static const char* block_statement_scope[] =
{
    NL, "kernel void block_statement_scope(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 0;"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  multiplier = 9;"
    NL, "  res[tid] = ^int(int num) { return multiplier * num; } (11) - 99;"
    NL, "}"
    NL
};

static const char* block_function_scope[] =
{
    NL, "int fnTest(int a)"
    NL, "{"
    NL, "  int localVar = 17;"
    NL, "  int (^functionBlock)(int) = ^(int num)"
    NL, "  {"
    NL, "    return localVar * num;"
    NL, "  };"
    NL, "  return 111 - functionBlock(a+1);"
    NL, "}"
    NL, "kernel void block_function_scope(__global int* res)"
    NL, "{"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  res[tid] = fnTest(5) - 9;"
    NL, "}"
    NL
};

static const char* block_nested_scope[] =
{
    NL, "kernel void block_nested_scope(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 3;"
    NL, "  int (^kernelBlock)(int) = ^(int num)"
    NL, "  {"
    NL, "    int (^innerBlock)(int) = ^(int n)"
    NL, "    {"
    NL, "      return multiplier * n;"
    NL, "    };"
    NL, "    return num * innerBlock(23);"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  multiplier = 8;"
    NL, "  res[tid] = kernelBlock(13) - 897;"
    NL, "}"
    NL
};

static const char* block_arg_struct[] =
{
    NL, "struct two_ints {"
    NL, "    short x;"
    NL, "    long y;"
    NL, "};"
    NL, "struct two_structs {"
    NL, "    struct two_ints a;"
    NL, "    struct two_ints b;"
    NL, "};"
    NL, "kernel void block_arg_struct(__global int* res)"
    NL, "{"
    NL, "  int (^kernelBlock)(struct two_ints, struct two_structs) = ^int(struct two_ints ti, struct two_structs ts)"
    NL, "  {"
    NL, "    return ti.x * ti.y * ts.a.x * ts.a.y * ts.b.x * ts.b.y;"
    NL, "  };"
    NL, "  struct two_ints i;"
    NL, "  i.x = 2;"
    NL, "  i.y = 3;"
    NL, "  struct two_structs s;"
    NL, "  s.a.x = 4;"
    NL, "  s.a.y = 5;"
    NL, "  s.b.x = 6;"
    NL, "  s.b.y = 7;"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  res[tid] = kernelBlock(i,s) - 5040;"
    NL, "}"
    NL
};

static const char* block_arg_types_mix[] =
{
    NL, "union number {"
    NL, "    long l;"
    NL, "    float f;"
    NL, "};"
    NL, "enum color {"
    NL, "    RED = 0,"
    NL, "    GREEN,"
    NL, "    BLUE" // Using this value - it is actualy "2"
    NL, "};"
    NL, "typedef int _INT ;"
    NL, "typedef char _ACHAR[3] ;"
    NL, "kernel void block_arg_types_mix(__global int* res)"
    NL, "{"
    NL, "  int (^kernelBlock)(_INT, _ACHAR, union number, enum color, int, int, int, int, int, int, int, int, int, int, int, int, int) ="
    NL, "    ^int(_INT bi, _ACHAR bch, union number bn, enum color bc, int i1, int i2, int i3, int i4, int i5, int i6, int i7, int i8,"
    NL, "      int i9, int i10, int i11, int i12, int i13)"
    NL, "  {"
    NL, "    return bi * bch[0] * bch[1] * bch[2] * bn.l * bc - i1 - i2 - i3 - i4 - i5 - i6 - i7 - i8 - i9 - i10 - i11 - i12 - i13;"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  _INT x = -5;"
    NL, "  _ACHAR char_arr = { 1, 2, 3 };"
    NL, "  union number n;"
    NL, "  n.l = 4;"
    NL, "  enum color c = BLUE;"
    NL, "  res[tid] = kernelBlock(x,char_arr,n,c,1,2,3,4,5,6,7,8,9,10,11,12,13) + 331;"
    NL, "}"
    NL
};

static const char* block_arg_pointer[] =
{
    NL, "struct two_ints {"
    NL, "    short x;"
    NL, "    long y;"
    NL, "};"
    NL, "kernel void block_arg_pointer(__global int* res)"
    NL, "{"
    NL, "  int (^kernelBlock)(struct two_ints*, struct two_ints*, int*, int*) = "
    NL, "    ^int(struct two_ints* bs1, struct two_ints* bs2, int* bi1, int* bi2)"
    NL, "  {"
    NL, "    return (*bs1).x * (*bs1).y * (*bs2).x * (*bs2).y * (*bi1) * (*bi2);"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  struct two_ints s[2];"
    NL, "  s[0].x = 4;"
    NL, "  s[0].y = 5;"
    NL, "  struct two_ints* ps = s + 1;"
    NL, "  (*ps).x = 6;"
    NL, "  (*ps).y = 7;"
    NL, "  int i = 2;"
    NL, "  int * pi = &i;"
    NL, "  res[tid] = kernelBlock(s,ps,&i,pi) - 3360;"
    NL, "}"
    NL
};

static const char* block_arg_global_p[] =
{
    NL, "kernel void block_arg_global_p(__global int* res)"
    NL, "{"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  typedef __global int* int_ptr_to_global_t;"
    NL, "  int_ptr_to_global_t (^kernelBlock)(__global int*, int) =^ int_ptr_to_global_t (__global int* bres, int btid)"
    NL, "  {"
    NL, "    bres[tid] = 5;"
    NL, "    return bres;"
    NL, "  };"
    NL, "  res = kernelBlock(res, tid);"
    NL, "  res[tid] -= 5;"
    NL, "}"
    NL
};

static const char* block_arg_const_p[] =
{
    NL, "constant int ci = 8;"
    NL, "kernel void block_arg_const_p(__global int* res)"
    NL, "{"
    NL, "  __constant int* (^kernelBlock)(__constant int*) = ^(__constant int* bpci)"
    NL, "  {"
    NL, "    return bpci;"
    NL, "  };"
    NL, "  constant int* pci = &ci;"
    NL, "  constant int* pci_check;"
    NL, "  pci_check = kernelBlock(pci);"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = pci == pci_check ? 0 : -1;"
    NL, "}"
    NL
};

static const char* block_ret_struct[] =
{
    NL, "kernel void block_ret_struct(__global int* res)"
    NL, "{"
    NL, "  struct A {"
    NL, "      int a;"
    NL, "  };      "
    NL, "  struct A (^kernelBlock)(struct A) = ^struct A(struct A a)"
    NL, "  {        "
    NL, "    a.a = 6;"
    NL, "    return a;"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  struct A aa;"
    NL, "  aa.a = 5;"
    NL, "  res[tid] = kernelBlock(aa).a - 6;"
    NL, "}"
    NL
};

static const char* block_arg_global_var[] =
{
    NL, "constant int gi = 8;"
    NL, "kernel void block_arg_global_var(__global int* res)"
    NL, "{"
    NL, "  int (^kernelBlock)(int) = ^(int bgi)"
    NL, "  {"
    NL, "    return bgi - 8;"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = kernelBlock(gi);"
    NL, "}"
    NL
};

static const char* block_in_for_init[] =
{
    NL, "kernel void block_in_for_init(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 3;"
    NL, "  int (^kernelBlock)(int) = ^(int num)"
    NL, "  {"
    NL, "    return num * multiplier;"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = 27;"
    NL, "  for(int i=kernelBlock(9); i>0; i--)"
    NL, "  {"
    NL, "       res[tid]--;"
    NL, "  }"
    NL, "}"
    NL
};

static const char* block_in_for_cond[] =
{
    NL, "kernel void block_in_for_cond(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 3;"
    NL, "  int (^kernelBlock)(int) = ^(int num)"
    NL, "  {"
    NL, "    return num * multiplier;"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = 39;"
    NL, "  for(int i=0; i<kernelBlock(13); i++)"
    NL, "  {"
    NL, "       res[tid]--;"
    NL, "  }"
    NL, "}"
    NL
};

static const char* block_in_for_iter[] =
{
    NL, "kernel void block_in_for_iter(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 2;"
    NL, "  int (^kernelBlock)(int) = ^(int num)"
    NL, "  {"
    NL, "    return num * multiplier;"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = 4;"
    NL, "  for(int i=2; i<17; i=kernelBlock(i))"
    NL, "  {"
    NL, "       res[tid]--;"
    NL, "  }"
    NL, "}"
    NL
};

static const char* block_in_while_cond[] =
{
    NL, "kernel void block_in_while_cond(__global int* res)"
    NL, "{"
    NL, "  int (^kernelBlock)(int) = ^(int num)"
    NL, "  {"
    NL, "    return res[num];"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = 27*(tid+1);"
    NL, "  while(kernelBlock(tid))"
    NL, "  {"
    NL, "      res[tid]--;"
    NL, "  }"
    NL, "}"
    NL
};

static const char* block_in_while_body[] =
{
    NL, "kernel void block_in_while_body(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 3;"
    NL, "  int (^kernelBlock)(int) = ^(int num)"
    NL, "  {"
    NL, "    return num * multiplier;"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  int i = 7;"
    NL, "  res[tid] = 3*(7+6+5+4+3+2+1);"
    NL, "  while(i)"
    NL, "  {"
    NL, "      res[tid]-=kernelBlock(i--);"
    NL, "  }"
    NL, "}"
    NL
};

static const char* block_in_do_while_body[] =
{
    NL, "kernel void block_in_do_while_body(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 3;"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  int i = 100;"
    NL, "  res[tid] = 3*5050;"
    NL, "  do"
    NL, "  {"
    NL, "      int (^kernelBlock)(int) = ^(int num)"
    NL, "      {"
    NL, "          return num * multiplier;"
    NL, "      };"
    NL, "      res[tid]-=kernelBlock(i--);"
    NL, "  } while(i);"
    NL, "}"
    NL
};

static const char* block_cond_statement[] =
{
    NL, "kernel void block_cond_statement(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 2;"
    NL, "  int (^kernelBlock)(int) = ^(int num)"
    NL, "  {"
    NL, "    return num * multiplier;"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = 120;"
    NL, "  res[tid] = (kernelBlock(2) == 4) ? res[tid] - 30              : res[tid] - 1;"
    NL, "  res[tid] = (kernelBlock(2) == 5) ? res[tid] - 3               : res[tid] - 30;"
    NL, "  res[tid] = (1)                   ? res[tid] - kernelBlock(15) : res[tid] - 7;"
    NL, "  res[tid] = (0)                   ? res[tid] - 13              : res[tid] - kernelBlock(15);"
    NL, "}"
    NL
};

static const char* block_in_if_cond[] =
{
    NL, "kernel void block_in_if_cond(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 2;"
    NL, "  int (^kernelBlock)(int) = ^(int num)"
    NL, "  {"
    NL, "    return num * multiplier;"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = 7;"
    NL, "  if (kernelBlock(5))"
    NL, "  {"
    NL, "      res[tid]-= 3;"
    NL, "  }"
    NL, "  if (kernelBlock(0))"
    NL, "  {"
    NL, "      res[tid]-= 2;"
    NL, "  }"
    NL, "  else"
    NL, "  {"
    NL, "      res[tid]-= 4;"
    NL, "  }"
    NL, "}"
    NL
};

static const char* block_in_if_branch[] =
{
    NL, "kernel void block_in_if_branch(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 2;"
    NL, "  int (^kernelBlock)(int) = ^(int num)"
    NL, "  {"
    NL, "    return num * multiplier;"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = 7;"
    NL, "  if (kernelBlock(5))"
    NL, "  {"
    NL, "      res[tid]-= ^(int num){ return num - 1; }(4);" // res[tid]-=3;
    NL, "  }"
    NL, "  if (kernelBlock(0))"
    NL, "  {"
    NL, "      res[tid]-= ^(int num){ return num - 1; }(3);" // res[tid]-=2;
    NL, "  }"
    NL, "  else"
    NL, "  {"
    NL, "      int (^ifBlock)(int) = ^(int num)"
    NL, "      {"
    NL, "          return num + 1;"
    NL, "      };"
    NL, "      res[tid]-= ifBlock(3);"                     // res[tid]-=4;
    NL, "  }"
    NL, "}"
    NL
};

static const char* block_switch_cond[] =
{
    NL, "kernel void block_switch_cond(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 2;"
    NL, "  int (^kernelBlock)(int) = ^(int num)"
    NL, "  {"
    NL, "    return num * multiplier;"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = 12;"
    NL, "  int i = 1;"
    NL, "  while(i <= 3)"
    NL, "  {"
    NL, "      switch (kernelBlock(i))"
    NL, "      {"
    NL, "          case 2:"
    NL, "              res[tid] = res[tid] - 2;"
    NL, "              break;"
    NL, "          case 4:"
    NL, "              res[tid] = res[tid] - 4;"
    NL, "              break;"
    NL, "          case 6:"
    NL, "              res[tid] = res[tid] - 6;"
    NL, "              break;"
    NL, "          default:"
    NL, "              break;"
    NL, "      }"
    NL, "      i++;"
    NL, "  }"
    NL, "}"
    NL
};

static const char* block_switch_case[] =
{
    NL, "kernel void block_switch_case(__global int* res)"
    NL, "{"
    NL, "  int multiplier = 2;"
    NL, "  int (^kernelBlock)(int) = ^(int num)"
    NL, "  {"
    NL, "    return num * multiplier;"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = 12;"
    NL, "  int i = 1;"
    NL, "  while(i <= 3)"
    NL, "  {"
    NL, "      switch (kernelBlock(i))"
    NL, "      {"
    NL, "          case 2:"
    NL, "              res[tid]-=^(int num){ return num - 1; }(3);" // res[tid]-=2;
    NL, "              break;"
    NL, "          case 4:"
    NL, "          {"
    NL, "              int (^caseBlock)(int) = ^(int num)"
    NL, "              {"
    NL, "                  return num + 1;"
    NL, "              };"
    NL, "              res[tid]-=caseBlock(3);"                   // res[tid]-=4;
    NL, "              break;"
    NL, "          }"
    NL, "          case 6:"
    NL, "              res[tid]-=kernelBlock(3);"                 // res[tid]-=6;
    NL, "              break;"
    NL, "          default:"
    NL, "              break;"
    NL, "      }"
    NL, "      i++;"
    NL, "  }"
    NL, "}"
    NL
};

// Accessing data from Block

static const char* block_access_program_data[] =
{
    NL, "int __constant globalVar1 = 7;"
    NL, "int __constant globalVar2 = 11;"
    NL, "int __constant globalVar3 = 13;"
    NL, "int (^__constant globalBlock)(int) = ^int(int num)"
    NL, "{"
    NL, "    return globalVar1 * num;"
    NL, "};"
    NL, "kernel void block_access_program_data(__global int* res)"
    NL, "{"
    NL, "    int (^ kernelBlock)(int) = ^int(int num)"
    NL, "    {"
    NL, "        return globalVar2 * num;"
    NL, "    };"
    NL, "    size_t tid = get_global_id(0);"
    NL, "    res[tid] = tid + 1;"
    NL, "    res[tid] = globalBlock(res[tid]);"
    NL, "    res[tid] = kernelBlock(res[tid]);"
    NL, "    res[tid] = ^(int num){ return globalVar3*num; }(res[tid]) - (7*11*13)*(tid + 1);"
    NL, "}"
    NL
};

static const char* block_access_kernel_data[] =
{
    NL, "kernel void block_access_kernel_data(__global int* res)"
    NL, "{"
    NL, "    int var1 = 7;"
    NL, "    int var2 = 11;"
    NL, "    int var3 = 13;"
    NL, "    int (^ kernelBlock)(int) = ^int(int num)"
    NL, "    {"
    NL, "        int (^ nestedBlock)(int) = ^int (int num)"
    NL, "        {"
    NL, "            return var1 * num;"
    NL, "        };"
    NL, "        return var2 * nestedBlock(num);"
    NL, "    };"
    NL, "    size_t tid = get_global_id(0);"
    NL, "    res[tid] = tid + 1;"
    NL, "    res[tid] = kernelBlock(res[tid]);"
    NL, "    res[tid] = ^(int num){ return var3*num; }(res[tid]) - (7*11*13)*(tid + 1);"
    NL, "}"
    NL
};

static const char* block_access_chained_data[] =
{
    NL, "kernel void block_access_chained_data(__global int* res)"
    NL, "{"
    NL, "    int (^ kernelBlock)(int) = ^int(int num)"
    NL, "    {"
    NL, "        int var1 = 7;"
    NL, "        int var2 = 11;"
    NL, "        int var3 = 13;"
    NL, "        int (^ nestedBlock1)(int) = ^int (int num)"
    NL, "        {"
    NL, "            int (^ nestedBlock2) (int) = ^int (int num)"
    NL, "            {"
    NL, "                return var2 * ^(int num){ return var3*num; }(num);"
    NL, "            };"
    NL, "            return var1 * nestedBlock2(num);"
    NL, "        };"
    NL, "        return nestedBlock1(num);"
    NL, "    };"
    NL, "    size_t tid = get_global_id(0);"
    NL, "    res[tid] = tid + 1;"
    NL, "    res[tid] = kernelBlock(res[tid]) - (7*11*13)*(tid + 1);"
    NL, "}"
    NL
};

static const char* block_access_volatile_data[] =
{
    NL, "kernel void block_access_volatile_data(__global int* res)"
    NL, "{"
    NL, "    int var1 = 7;"
    NL, "    int var2 = 11;"
    NL, "    volatile int var3 = 13;"
    NL, ""
    NL, "    int (^ kernelBlock)(int) = ^int(int num)"
    NL, "    {"
    NL, "        int (^ nestedBlock)(int) = ^int (int num)"
    NL, "        {"
    NL, "            return var1 * num;"
    NL, "        };"
    NL, "        return var2 * nestedBlock(num);"
    NL, "    };"
    NL, "    size_t tid = get_global_id(0);"
    NL, "    res[tid] = tid + 1;"
    NL, "    res[tid] = kernelBlock(res[tid]);"
    NL, "    res[tid] = ^(int num){ return var3*num; }(res[tid]) - (7*11*13)*(tid + 1);"
    NL, "}"
    NL
};

static const char* block_typedef_kernel[] =
{
    NL, "kernel void block_typedef_kernel(__global int* res)"
    NL, "{"
    NL, "  typedef int* (^block_t)(int*);"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  int i[4] = { 3, 4, 4, 1 };"
    NL, "  int *temp = i; // workaround clang bug"
    NL, "  block_t kernelBlock = ^(int* pi)"
    NL, "  {"
    NL, "    block_t b = ^(int* n) { return n - 1; };"
    NL, "    return pi + *(b(temp+4));"
    NL, "  };"
    NL, "  switch (*(kernelBlock(i))) {"
    NL, "    case 4:"
    NL, "      res[tid] += *(kernelBlock(i+1));"
    NL, "      break;"
    NL, "    default:"
    NL, "      res[tid] = -100;"
    NL, "      break;"
    NL, "  }"
    NL, "  res[tid] += *(kernelBlock(i)) - 7;"
    NL, "}"
    NL
};

static const char* block_typedef_func[] =
{
    NL, "int func(int fi)"
    NL, "{"
    NL, "  typedef int (^block_t)(int);"
    NL, "  const block_t funcBlock = ^(int bi)"
    NL, "  {"
    NL, "    typedef short (^block2_t)(short);"
    NL, "    block2_t nestedBlock = ^(short ni)"
    NL, "    {"
    NL, "      return (short)(ni - 1);"
    NL, "    };"
    NL, "    return bi * nestedBlock(3);"
    NL, "  };"
    NL, "  return funcBlock(fi * 2);"
    NL, "}"
    NL, "kernel void block_typedef_func(__global int* res)"
    NL, "{"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  res[tid] = func(1) - 4;"
    NL, "}"
    NL
};

static const char* block_typedef_stmnt_if[] =
{
    NL, "kernel void block_typedef_stmnt_if(__global int* res)"
    NL, "{      "
    NL, "  int flag = 1;"
    NL, "  int sum = 0;"
    NL, "  if (flag) {"
    NL, "    typedef int (^block_t)(int);"
    NL, "    const block_t kernelBlock = ^(int bi)"
    NL, "    {"
    NL, "      block_t b = ^(int bi)"
    NL, "      {"
    NL, "        return bi + 1;"
    NL, "      };"
    NL, "      return bi + b(1);"
    NL, "    };"
    NL, "    sum = kernelBlock(sum);"
    NL, "  }"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = sum - 2;"
    NL, "}"
    NL
};

static const char* block_typedef_loop[] =
{
    NL, "kernel void block_typedef_loop(__global int* res)"
    NL, "{      "
    NL, "  int sum = -1;"
    NL, "  for (int i = 0; i < 3; i++) {"
    NL, "    typedef int (^block_t)(void);"
    NL, "    const block_t kernelBlock = ^()"
    NL, "    {"
    NL, "      return i + 1;"
    NL, "    };"
    NL, "    sum += kernelBlock();"
    NL, "  }"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = sum - 5;"
    NL, "}"
    NL
};

static const char* block_typedef_mltpl_func[] =
{
    NL, "int func(int fi)"
    NL, "{"
    NL, "  typedef int (^block_t)(int);"
    NL, "  typedef int (^block2_t)(int);"
    NL, "  const block_t funcBlock1 = ^(int bi) { return bi; };"
    NL, "  const block2_t funcBlock2 = ^(int bi)"
    NL, "  {"
    NL, "    typedef short (^block3_t)(short);"
    NL, "    typedef short (^block4_t)(short);"
    NL, "    const block3_t nestedBlock1 = ^(short ni)"
    NL, "    {"
    NL, "      return (short)(ni - 1);"
    NL, "    };"
    NL, "    const block4_t nestedBlock2 = ^(short ni)"
    NL, "    {"
    NL, "      return (short)(ni - 2);"
    NL, "    };"
    NL, "    return bi * nestedBlock1(3) * nestedBlock2(3);"
    NL, "  };"
    NL, "  return funcBlock2(fi * 2) + funcBlock1(1);"
    NL, "}"
    NL, "kernel void block_typedef_mltpl_func(__global int* res)"
    NL, "{"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  typedef int (^block1_t)(int);"
    NL, "  typedef int (^block2_t)(int);"
    NL, "  const block1_t kernelBlock1 = ^(int bi) { return bi + 8; };"
    NL, "  const block2_t kernelBlock2 = ^(int bi) { return bi + 3; };"
    NL, "  res[tid] = func(1) -  kernelBlock1(2) / kernelBlock2(-1);"
    NL, "}"
    NL
};

static const char* block_typedef_mltpl_stmnt[] =
{
    NL, "kernel void block_typedef_mltpl_stmnt(__global int* res)"
    NL, "{"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  int a;"
    NL, "  do"
    NL, "  {"
    NL, "    typedef float (^blockf_t)(float);"
    NL, "    typedef int (^blocki_t)(int);"
    NL, "    const blockf_t blockF = ^(float bi) { return (float)(bi + 3.3); };"
    NL, "    const blocki_t blockI = ^(int bi) { return bi + 2; };"
    NL, "    if ((blockF(.0)-blockI(0)) > 0)"
    NL, "    {"
    NL, "      typedef uint (^block_t)(uint);"
    NL, "      const block_t nestedBlock = ^(uint bi) { return (uint)(bi + 4); };"
    NL, "      a = nestedBlock(1) + nestedBlock(2);"
    NL, "      break;"
    NL, "    }"
    NL, "  } while(1);  "
    NL, "  res[tid] = a - 11;"
    NL, "}"
    NL
};

static const char* block_typedef_mltpl_g[] =
{
    NL, "typedef int (^block1_t)(float, int); "
    NL, "constant block1_t b1 = ^(float fi, int ii) { return (int)(ii + fi); };"
    NL, "typedef int (^block2_t)(float, int);"
    NL, "constant block2_t b2 = ^(float fi, int ii) { return (int)(ii + fi); };"
    NL, "typedef float (^block3_t)(int, int);"
    NL, "constant block3_t b3 = ^(int i1, int i2) { return (float)(i1 + i2); };"
    NL, "typedef int (^block4_t)(float, float);"
    NL, "kernel void block_typedef_mltpl_g(__global int* res)"
    NL, "{"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  block4_t b4 = ^(float f1, float f2) { return (int)(f1 + f2); };"
    NL, "  res[tid] = b1(1.1, b2(1.1, 1)) - b4(b3(1,1), 1.1);"
    NL, "}"
    NL
};

static const char* block_literal[] =
{
    NL, "int func()"
    NL, "{"
    NL, "  return ^(int i) {"
    NL, "    return ^(ushort us)"
    NL, "    {"
    NL, "      return (int)us + i;"
    NL, "    }(3);"
    NL, "  }(7) - 10;"
    NL, "}"
    NL, "kernel void block_literal(__global int* res)"
    NL, "{"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  res[tid] = func();"
    NL, "}"
    NL
};

static const char* block_complex[] =
{
    NL, "kernel void block_complex(__global int* res)"
    NL, "{"
    NL, "  int (^kernelBlock)(int) = ^(int num)"
    NL, "  {"
    NL, "    int result = 1;"
    NL, "    for (int i = 0; i < num; i++)"
    NL, "    {"
    NL, "      switch(i)"
    NL, "      {"
    NL, "      case 0:"
    NL, "      case 1:"
    NL, "      case 2:"
    NL, "        result += i;"
    NL, "        break;"
    NL, "      case 3:"
    NL, "        if (result < num)"
    NL, "          result += i;"
    NL, "        else"
    NL, "          result += i * 2;"
    NL, "        break;"
    NL, "      case 4:"
    NL, "        while (1)"
    NL, "        {"
    NL, "          result++;"
    NL, "          if (result)"
    NL, "            goto ret;"
    NL, "        }"
    NL, "        break;"
    NL, "      default:"
    NL, "        return 777;"
    NL, "      }"
    NL, "    }"
    NL, "    ret: ;"
    NL, "    while (num) {"
    NL, "      num--;"
    NL, "      if (num % 2 == 0)"
    NL, "        continue;"
    NL, "      result++;"
    NL, "    }"
    NL, "    return result;"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  res[tid] = kernelBlock(7) - 11;"
    NL, "}"
    NL
};

static const char* block_empty[] =
{
    NL, "kernel void block_empty(__global int* res)"
    NL, "{"
    NL, "  void (^kernelBlock)(void) = ^(){};"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  kernelBlock();"
    NL, "  res[tid] = 0;"
    NL, "}"
    NL
};

static const char* block_builtin[] =
{
    NL, "kernel void block_builtin(__global int* res)"
    NL, "{"
    NL, "  int b = 3;"
    NL, "  int (^kernelBlock)(int) = ^(int a)"
    NL, "  {"
    NL, "    return (int)abs(a - b);"
    NL, "  };"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  res[tid] = -1;"
    NL, "  res[tid] = kernelBlock(2) - 1;"
    NL, "}"
    NL
};

static const char* block_barrier[] =
{
    NL, "kernel void block_barrier(__global int* res)"
    NL, "{"
    NL, "  int b = 3;"
    NL, "  size_t tid = get_global_id(0);"
    NL, "  size_t lsz = get_local_size(0);"
    NL, "  size_t gid = get_group_id(0);"
    NL, "  size_t idx = gid*lsz;"
    NL, ""
    NL, "  res[tid]=lsz;"
    NL, "  barrier(CLK_GLOBAL_MEM_FENCE);"
    NL, "  int (^kernelBlock)(int) = ^(int a)"
    NL, "  {"
    NL, "    atomic_dec(res+idx);"
    NL, "    barrier(CLK_GLOBAL_MEM_FENCE);"
    NL, "    return (int)abs(a - b) - (res[idx] != 0 ? 0 : 1);"
    NL, "  };"
    NL, ""
    NL, "  int d = kernelBlock(2);"
    NL, "  res[tid] = d;"
    NL, "}"
    NL
};



static const kernel_src sources_execute_block[] =
{
    // Simple blocks
    KERNEL(block_global_scope),
    KERNEL(block_kernel_scope),
    KERNEL(block_statement_scope),
    KERNEL(block_function_scope),
    KERNEL(block_nested_scope),

    // Kernels with Block in for/while/if/switch
    KERNEL(block_in_for_init),
    KERNEL(block_in_for_cond),
    KERNEL(block_in_for_iter),
    KERNEL(block_in_while_cond),
    KERNEL(block_in_while_body),
    KERNEL(block_in_do_while_body),
    KERNEL(block_cond_statement),
    KERNEL(block_in_if_cond),
    KERNEL(block_in_if_branch),
    KERNEL(block_switch_cond),
    KERNEL(block_switch_case),
    KERNEL(block_literal),

    // Accessing data from block
    KERNEL(block_access_program_data),
    KERNEL(block_access_kernel_data),
    KERNEL(block_access_chained_data),
    KERNEL(block_access_volatile_data),

    // Block args
    KERNEL(block_arg_struct),
    KERNEL(block_arg_types_mix),
    KERNEL(block_arg_pointer),
    KERNEL(block_arg_global_p),
    KERNEL(block_arg_const_p),
    KERNEL(block_ret_struct),
    KERNEL(block_arg_global_var),

    // Block in typedef
    KERNEL(block_typedef_kernel),
    KERNEL(block_typedef_func),
    KERNEL(block_typedef_stmnt_if),
    KERNEL(block_typedef_loop),
    KERNEL(block_typedef_mltpl_func),
    KERNEL(block_typedef_mltpl_stmnt),
    KERNEL(block_typedef_mltpl_g),

    // Non - trivial blocks
    KERNEL(block_complex),
    KERNEL(block_empty),
    KERNEL(block_builtin),
    KERNEL(block_barrier),

};
static const size_t num_kernels_execute_block = arr_size(sources_execute_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_execute_block(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
{
    size_t i;
    size_t ret_len;
    cl_int n, err_ret, res = 0;
    clCommandQueueWrapper dev_queue;
    cl_int kernel_results[MAX_GWS] = {0xDEADBEEF};

    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");

    size_t global_size = MAX_GWS;
    size_t local_size = (max_local_size > global_size/16) ? global_size/16 : max_local_size;

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

        log_info("Running '%s' kernel (%d of %d) ...\n", sources_execute_block[i].kernel_name, i + 1, num_kernels_execute_block);
        err_ret = run_n_kernel_args(context, queue, sources_execute_block[i].lines, sources_execute_block[i].num_lines, sources_execute_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_execute_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_execute_block[i].kernel_name, n, kernel_results[n])) { ++failCnt; res = -1; }
        else log_info("'%s' kernel is OK.\n", sources_execute_block[i].kernel_name);
    }

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

    return res;
}


#endif