ceph/src/boost/libs/compute/test/test_kernel.cpp

   1 //---------------------------------------------------------------------------//
   2 // Copyright (c) 2013 Kyle Lutz <kyle.r.lutz@gmail.com>
   3 //
   4 // Distributed under the Boost Software License, Version 1.0
   5 // See accompanying file LICENSE_1_0.txt or copy at
   6 // http://www.boost.org/LICENSE_1_0.txt
   7 //
   8 // See http://boostorg.github.com/compute for more information.
   9 //---------------------------------------------------------------------------//
  10
  11 #define BOOST_TEST_MODULE TestKernel
  12 #include <boost/test/unit_test.hpp>
  13
  14 #include <boost/compute/buffer.hpp>
  15 #include <boost/compute/kernel.hpp>
  16 #include <boost/compute/system.hpp>
  17 #include <boost/compute/utility/source.hpp>
  18
  19 #include "context_setup.hpp"
  20 #include "check_macros.hpp"
  21
  22 namespace compute = boost::compute;
  23
  24 BOOST_AUTO_TEST_CASE(name)
  25 {
  26     compute::kernel foo = compute::kernel::create_with_source(
  27         "__kernel void foo(int x) { }", "foo", context
  28     );
  29     BOOST_CHECK_EQUAL(foo.name(), "foo");
  30
  31     compute::kernel bar = compute::kernel::create_with_source(
  32         "__kernel void bar(float x) { }", "bar", context
  33     );
  34     BOOST_CHECK_EQUAL(bar.name(), "bar");
  35 }
  36
  37 BOOST_AUTO_TEST_CASE(arity)
  38 {
  39     compute::kernel foo = compute::kernel::create_with_source(
  40         "__kernel void foo(int x) { }", "foo", context
  41     );
  42     BOOST_CHECK_EQUAL(foo.arity(), size_t(1));
  43
  44     compute::kernel bar = compute::kernel::create_with_source(
  45         "__kernel void bar(float x, float y) { }", "bar", context
  46     );
  47     BOOST_CHECK_EQUAL(bar.arity(), size_t(2));
  48
  49     compute::kernel baz = compute::kernel::create_with_source(
  50         "__kernel void baz(char x, char y, char z) { }", "baz", context
  51     );
  52     BOOST_CHECK_EQUAL(baz.arity(), size_t(3));
  53 }
  54
  55 BOOST_AUTO_TEST_CASE(set_buffer_arg)
  56 {
  57     const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
  58         __kernel void foo(__global int *x, __global int *y)
  59         {
  60             x[get_global_id(0)] = -y[get_global_id(0)];
  61         }
  62     );
  63
  64     compute::kernel foo =
  65         compute::kernel::create_with_source(source, "foo", context);
  66
  67     compute::buffer x(context, 16);
  68     compute::buffer y(context, 16);
  69
  70     foo.set_arg(0, x);
  71     foo.set_arg(1, y.get());
  72 }
  73
  74 BOOST_AUTO_TEST_CASE(get_work_group_info)
  75 {
  76     const char source[] =
  77     "__kernel void sum(__global const float *input,\n"
  78     "                  __global float *output)\n"
  79     "{\n"
  80     "    __local float scratch[16];\n"
  81     "    const uint gid = get_global_id(0);\n"
  82     "    const uint lid = get_local_id(0);\n"
  83     "    if(lid < 16)\n"
  84     "        scratch[lid] = input[gid];\n"
  85     "}\n";
  86
  87     compute::program program =
  88         compute::program::create_with_source(source, context);
  89
  90     program.build();
  91
  92     compute::kernel kernel = program.create_kernel("sum");
  93
  94     using compute::ulong_;
  95
  96     // get local memory size
  97     kernel.get_work_group_info<ulong_>(device, CL_KERNEL_LOCAL_MEM_SIZE);
  98
  99     // check work group size
 100     size_t work_group_size =
 101         kernel.get_work_group_info<size_t>(device, CL_KERNEL_WORK_GROUP_SIZE);
 102     BOOST_CHECK(work_group_size >= 1);
 103 }
 104
 105 #ifndef BOOST_COMPUTE_NO_VARIADIC_TEMPLATES
 106 BOOST_AUTO_TEST_CASE(kernel_set_args)
 107 {
 108     compute::kernel k = compute::kernel::create_with_source(
 109         "__kernel void test(int x, float y, char z) { }", "test", context
 110     );
 111
 112     k.set_args(4, 2.4f, 'a');
 113 }
 114 #endif // BOOST_COMPUTE_NO_VARIADIC_TEMPLATES
 115
 116 #ifdef BOOST_COMPUTE_CL_VERSION_1_2
 117 BOOST_AUTO_TEST_CASE(get_arg_info)
 118 {
 119     REQUIRES_OPENCL_VERSION(1, 2);
 120
 121     const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
 122         __kernel void sum_kernel(__global const int *input,
 123                                  const uint size,
 124                                  __global int *result)
 125         {
 126             int sum = 0;
 127             for(uint i = 0; i < size; i++){
 128                 sum += input[i];
 129             }
 130             *result = sum;
 131         }
 132     );
 133
 134     compute::program program =
 135         compute::program::create_with_source(source, context);
 136
 137     program.build("-cl-kernel-arg-info");
 138
 139     compute::kernel kernel = program.create_kernel("sum_kernel");
 140
 141     BOOST_CHECK_EQUAL(kernel.get_info<CL_KERNEL_NUM_ARGS>(), compute::uint_(3));
 142
 143     BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(0, CL_KERNEL_ARG_TYPE_NAME), "int*");
 144     BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(0, CL_KERNEL_ARG_NAME), "input");
 145     BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(1, CL_KERNEL_ARG_TYPE_NAME), "uint");
 146     BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(1, CL_KERNEL_ARG_NAME), "size");
 147     BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(2, CL_KERNEL_ARG_TYPE_NAME), "int*");
 148     BOOST_CHECK_EQUAL(kernel.get_arg_info<std::string>(2, CL_KERNEL_ARG_NAME), "result");
 149
 150     BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_TYPE_NAME>(0), "int*");
 151     BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_NAME>(0), "input");
 152     BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_TYPE_NAME>(1), "uint");
 153     BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_NAME>(1), "size");
 154     BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_TYPE_NAME>(2), "int*");
 155     BOOST_CHECK_EQUAL(kernel.get_arg_info<CL_KERNEL_ARG_NAME>(2), "result");
 156 }
 157 #endif // BOOST_COMPUTE_CL_VERSION_1_2
 158
 159 #ifdef BOOST_COMPUTE_CL_VERSION_2_0
 160 #ifndef CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE_KHR
 161     #define CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE_KHR CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE
 162 #endif
 163 #ifndef CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE_KHR
 164     #define CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE_KHR CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE
 165 #endif
 166 BOOST_AUTO_TEST_CASE(get_sub_group_info_ext)
 167 {
 168     compute::kernel k = compute::kernel::create_with_source(
 169         "__kernel void test(float x) { }", "test", context
 170     );
 171
 172     // get_sub_group_info(const device&, cl_kernel_sub_group_info, const std::vector<size_t>)
 173     std::vector<size_t> local_work_size(2, size_t(64));
 174     boost::optional<size_t> count = k.get_sub_group_info<size_t>(
 175         device,
 176         CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE_KHR,
 177         local_work_size
 178     );
 179
 180     if(device.check_version(2, 1))
 181     {
 182         BOOST_CHECK(count);
 183     }
 184     else if(device.check_version(2, 0) && device.supports_extension("cl_khr_subgroups"))
 185     {
 186         // for device with cl_khr_subgroups it should return some value
 187         BOOST_CHECK(count);
 188     }
 189     else
 190     {
 191         // for device without cl_khr_subgroups ext it should return null optional
 192         BOOST_CHECK(count == boost::none);
 193     }
 194
 195     // get_sub_group_info(const device&, cl_kernel_sub_group_info, const size_t, const void *)
 196     count = k.get_sub_group_info<size_t>(
 197         device,
 198         CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE_KHR,
 199         2 * sizeof(size_t),
 200         &local_work_size[0]
 201     );
 202
 203     if(device.check_version(2, 1))
 204     {
 205         BOOST_CHECK(count);
 206     }
 207     else if(device.check_version(2, 0) && device.supports_extension("cl_khr_subgroups"))
 208     {
 209         // for device with cl_khr_subgroups it should return some value
 210         BOOST_CHECK(count);
 211     }
 212     else
 213     {
 214         // for device without cl_khr_subgroups ext it should return null optional
 215         BOOST_CHECK(count == boost::none);
 216     }
 217 }
 218 #endif // BOOST_COMPUTE_CL_VERSION_2_0
 219
 220 #ifdef BOOST_COMPUTE_CL_VERSION_2_1
 221 BOOST_AUTO_TEST_CASE(get_sub_group_info_core)
 222 {
 223     compute::kernel k = compute::kernel::create_with_source(
 224         "__kernel void test(float x) { }", "test", context
 225     );
 226
 227     // get_sub_group_info(const device&, cl_kernel_sub_group_info, const size_t)
 228     boost::optional<std::vector<size_t>> local_size =
 229         k.get_sub_group_info<std::vector<size_t> >(
 230             device,
 231             CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT,
 232             size_t(1)
 233         );
 234
 235     if(device.check_version(2, 1))
 236     {
 237         // for 2.1 devices it should return some value
 238         BOOST_CHECK(local_size);
 239         BOOST_CHECK(local_size.value().size() == 3);
 240     }
 241     else
 242     {
 243         // for 1.x and 2.0 devices it should return null optional,
 244         // because CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT is not
 245         // supported by cl_khr_subgroups (2.0 ext)
 246         BOOST_CHECK(local_size == boost::none);
 247     }
 248
 249     // get_sub_group_info(const device&, cl_kernel_sub_group_info, const size_t)
 250     boost::optional<size_t> local_size_simple =
 251         k.get_sub_group_info<size_t>(
 252             device,
 253             CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT,
 254             size_t(1)
 255         );
 256
 257     if(device.check_version(2, 1))
 258     {
 259         // for 2.1 devices it should return some value
 260         BOOST_CHECK(local_size_simple);
 261     }
 262     else
 263     {
 264         // for 1.x and 2.0 devices it should return null optional,
 265         // because CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT is not
 266         // supported by cl_khr_subgroups (2.0 ext)
 267         BOOST_CHECK(local_size_simple == boost::none);
 268     }
 269
 270     // get_sub_group_info(const device&, cl_kernel_sub_group_info)
 271     boost::optional<size_t> max =
 272         k.get_sub_group_info<size_t>(
 273             device,
 274             CL_KERNEL_MAX_NUM_SUB_GROUPS
 275         );
 276
 277     if(device.check_version(2, 1))
 278     {
 279         // for 2.1 devices it should return some value
 280         BOOST_CHECK(max);
 281     }
 282     else
 283     {
 284         // for 1.x and 2.0 devices it should return null optional,
 285         // because CL_KERNEL_MAX_NUM_SUB_GROUPS is not
 286         // supported by cl_khr_subgroups (2.0 ext)
 287         BOOST_CHECK(max == boost::none);
 288     }
 289 }
 290 #endif // BOOST_COMPUTE_CL_VERSION_2_1
 291
 292 #ifdef BOOST_COMPUTE_CL_VERSION_2_1
 293 BOOST_AUTO_TEST_CASE(clone_kernel)
 294 {
 295     REQUIRES_OPENCL_PLATFORM_VERSION(2, 1);
 296
 297     compute::kernel k1 = compute::kernel::create_with_source(
 298         "__kernel void test(__global int * x) { x[get_global_id(0)] = get_global_id(0); }",
 299         "test", context
 300     );
 301
 302     compute::buffer x(context, 5 * sizeof(compute::int_));
 303     k1.set_arg(0, x);
 304
 305     // Clone k1 kernel
 306     compute::kernel k2 = k1.clone();
 307     // After clone k2 0th argument (__global float * x) should be set,
 308     // so we should be able to enqueue k2 kernel without problems
 309     queue.enqueue_1d_range_kernel(k2, 0, x.size() / sizeof(compute::int_), 0).wait();
 310 }
 311 #endif // BOOST_COMPUTE_CL_VERSION_2_1
 312
 313 BOOST_AUTO_TEST_SUITE_END()