ceph/src/boost/libs/compute/test/test_program.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 TestProgram
  12 #include <boost/test/unit_test.hpp>
  13
  14 // disable the automatic kernel compilation debug messages. this allows the
  15 // test for program to check that compilation error exceptions are properly
  16 // thrown when invalid kernel code is passed to program::build().
  17 #undef BOOST_COMPUTE_DEBUG_KERNEL_COMPILATION
  18
  19 #include <boost/compute/exception/program_build_failure.hpp>
  20 #include <boost/compute/kernel.hpp>
  21 #include <boost/compute/system.hpp>
  22 #include <boost/compute/program.hpp>
  23 #include <boost/compute/utility/source.hpp>
  24
  25 #include "quirks.hpp"
  26 #include "context_setup.hpp"
  27
  28 namespace compute = boost::compute;
  29
  30 const char source[] =
  31     "__kernel void foo(__global float *x, const uint n) { }\n"
  32     "__kernel void bar(__global int *x, __global int *y) { }\n";
  33
  34
  35 BOOST_AUTO_TEST_CASE(get_program_info)
  36 {
  37     // create program
  38     boost::compute::program program =
  39         boost::compute::program::create_with_source(source, context);
  40
  41     // build program
  42     program.build();
  43
  44     // check program info
  45 #ifndef BOOST_COMPUTE_USE_OFFLINE_CACHE
  46     BOOST_CHECK(program.source().empty() == false);
  47 #endif
  48     BOOST_CHECK(program.get_context() == context);
  49 }
  50
  51 BOOST_AUTO_TEST_CASE(program_source)
  52 {
  53     // create program from source
  54     boost::compute::program program =
  55         boost::compute::program::create_with_source(source, context);
  56
  57     BOOST_CHECK_EQUAL(std::string(source), program.source());
  58 }
  59
  60 BOOST_AUTO_TEST_CASE(program_multiple_sources)
  61 {
  62     std::vector<std::string> sources;
  63     sources.push_back("__kernel void foo(__global int* x) { }\n");
  64     sources.push_back("__kernel void bar(__global float* y) { }\n");
  65
  66     // create program from sources
  67     boost::compute::program program =
  68         boost::compute::program::create_with_source(sources, context);
  69     program.build();
  70
  71     boost::compute::kernel foo = program.create_kernel("foo");
  72     boost::compute::kernel bar = program.create_kernel("bar");
  73 }
  74
  75 BOOST_AUTO_TEST_CASE(program_source_no_file)
  76 {
  77     // create program from a non-existant source file
  78     // and verifies it throws.
  79     BOOST_CHECK_THROW(boost::compute::program program =
  80                       boost::compute::program::create_with_source_file
  81                       (std::string(), context),
  82                       std::ios_base::failure);
  83 }
  84
  85 BOOST_AUTO_TEST_CASE(create_kernel)
  86 {
  87     boost::compute::program program =
  88         boost::compute::program::create_with_source(source, context);
  89     program.build();
  90
  91     boost::compute::kernel foo = program.create_kernel("foo");
  92     boost::compute::kernel bar = program.create_kernel("bar");
  93
  94     // try to create a kernel that doesn't exist
  95     BOOST_CHECK_THROW(program.create_kernel("baz"), boost::compute::opencl_error);
  96 }
  97
  98 BOOST_AUTO_TEST_CASE(create_with_binary)
  99 {
 100     // create program from source
 101     boost::compute::program source_program =
 102         boost::compute::program::create_with_source(source, context);
 103     source_program.build();
 104
 105     // create kernels in source program
 106     boost::compute::kernel source_foo_kernel = source_program.create_kernel("foo");
 107     boost::compute::kernel source_bar_kernel = source_program.create_kernel("bar");
 108
 109     // check source kernels
 110     BOOST_CHECK_EQUAL(source_foo_kernel.name(), std::string("foo"));
 111     BOOST_CHECK_EQUAL(source_bar_kernel.name(), std::string("bar"));
 112
 113     // get binary
 114     std::vector<unsigned char> binary = source_program.binary();
 115
 116     // create program from binary
 117     boost::compute::program binary_program =
 118         boost::compute::program::create_with_binary(binary, context);
 119     binary_program.build();
 120
 121     // create kernels in binary program
 122     boost::compute::kernel binary_foo_kernel = binary_program.create_kernel("foo");
 123     boost::compute::kernel binary_bar_kernel = binary_program.create_kernel("bar");
 124
 125     // check binary kernels
 126     BOOST_CHECK_EQUAL(binary_foo_kernel.name(), std::string("foo"));
 127     BOOST_CHECK_EQUAL(binary_bar_kernel.name(), std::string("bar"));
 128 }
 129
 130 #ifdef BOOST_COMPUTE_CL_VERSION_2_1
 131 BOOST_AUTO_TEST_CASE(create_with_il)
 132 {
 133     REQUIRES_OPENCL_VERSION(2, 1);
 134
 135     size_t device_address_space_size = device.address_bits();
 136     std::string file_path(BOOST_COMPUTE_TEST_DATA_PATH);
 137     if(device_address_space_size == 64)
 138     {
 139         file_path += "/program.spirv64";
 140     }
 141     else
 142     {
 143         file_path += "/program.spirv32";
 144     }
 145
 146     // create program from il
 147     boost::compute::program il_program;
 148     BOOST_CHECK_NO_THROW(
 149         il_program = boost::compute::program::create_with_il_file(
 150             file_path, context
 151         )
 152     );
 153     BOOST_CHECK_NO_THROW(il_program.build());
 154
 155     // create kernel (to check if program was loaded correctly)
 156     BOOST_CHECK_NO_THROW(il_program.create_kernel("foobar"));
 157 }
 158
 159 BOOST_AUTO_TEST_CASE(get_program_il_binary)
 160 {
 161     REQUIRES_OPENCL_VERSION(2, 1);
 162
 163     size_t device_address_space_size = device.address_bits();
 164     std::string file_path(BOOST_COMPUTE_TEST_DATA_PATH);
 165     if(device_address_space_size == 64)
 166     {
 167         file_path += "/program.spirv64";
 168     }
 169     else
 170     {
 171         file_path += "/program.spirv32";
 172     }
 173
 174     // create program from il
 175     boost::compute::program il_program;
 176     BOOST_CHECK_NO_THROW(
 177         il_program = boost::compute::program::create_with_il_file(
 178             file_path, context
 179         )
 180     );
 181     BOOST_CHECK_NO_THROW(il_program.build());
 182
 183     std::vector<unsigned char> il_binary;
 184     BOOST_CHECK_NO_THROW(il_binary = il_program.il_binary());
 185
 186     // create program from loaded il binary
 187     BOOST_CHECK_NO_THROW(
 188         il_program = boost::compute::program::create_with_il(il_binary, context)
 189     );
 190     BOOST_CHECK_NO_THROW(il_program.build());
 191
 192     // create kernel (to check if program was loaded correctly)
 193     BOOST_CHECK_NO_THROW(il_program.create_kernel("foobar"));
 194 }
 195
 196 BOOST_AUTO_TEST_CASE(get_program_il_binary_empty)
 197 {
 198     REQUIRES_OPENCL_VERSION(2, 1);
 199
 200     boost::compute::program program;
 201     BOOST_CHECK_NO_THROW(
 202         program = boost::compute::program::create_with_source(source, context)
 203     );
 204     BOOST_CHECK_NO_THROW(program.build());
 205
 206     std::vector<unsigned char> il_binary;
 207     il_binary = program.il_binary();
 208     BOOST_CHECK(il_binary.empty());
 209 }
 210 #endif // BOOST_COMPUTE_CL_VERSION_2_1
 211
 212 BOOST_AUTO_TEST_CASE(create_with_source_doctest)
 213 {
 214 //! [create_with_source]
 215 std::string source = "__kernel void foo(__global int *data) { }";
 216
 217 boost::compute::program foo_program =
 218     boost::compute::program::create_with_source(source, context);
 219 //! [create_with_source]
 220
 221     foo_program.build();
 222 }
 223
 224 #ifdef BOOST_COMPUTE_CL_VERSION_1_2
 225 BOOST_AUTO_TEST_CASE(compile_and_link)
 226 {
 227     REQUIRES_OPENCL_VERSION(1,2);
 228
 229     if(!supports_compile_program(device) || !supports_link_program(device)) {
 230         return;
 231     }
 232
 233     // create the library program
 234     const char library_source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
 235         // for some reason the apple opencl compilers complains if a prototype
 236         // for the square() function is not available, so we add it here
 237         T square(T);
 238
 239         // generic square function definition
 240         T square(T x) { return x * x; }
 241     );
 242
 243     compute::program library_program =
 244         compute::program::create_with_source(library_source, context);
 245
 246     library_program.compile("-DT=int");
 247
 248     // create the kernel program
 249     const char kernel_source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
 250         // forward declare square function
 251         extern int square(int);
 252
 253         // square kernel definition
 254         __kernel void square_kernel(__global int *x)
 255         {
 256             x[0] = square(x[0]);
 257         }
 258     );
 259
 260     compute::program square_program =
 261         compute::program::create_with_source(kernel_source, context);
 262
 263     square_program.compile();
 264
 265     // link the programs
 266     std::vector<compute::program> programs;
 267     programs.push_back(library_program);
 268     programs.push_back(square_program);
 269
 270     compute::program linked_program =
 271         compute::program::link(programs, context);
 272
 273     // create the square kernel
 274     compute::kernel square_kernel =
 275         linked_program.create_kernel("square_kernel");
 276     BOOST_CHECK_EQUAL(square_kernel.name(), "square_kernel");
 277 }
 278
 279 BOOST_AUTO_TEST_CASE(compile_and_link_with_headers)
 280 {
 281     REQUIRES_OPENCL_VERSION(1,2);
 282
 283     if(!supports_compile_program(device) || !supports_link_program(device)) {
 284         return;
 285     }
 286
 287     // create the header programs
 288     const char square_header_source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
 289         T square(T x) { return x * x; }
 290     );
 291     const char div2_header_source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
 292         T div2(T x) { return x / 2; }
 293     );
 294
 295     compute::program square_header_program =
 296         compute::program::create_with_source(square_header_source, context);
 297     compute::program div2_header_program =
 298         compute::program::create_with_source(div2_header_source, context);
 299
 300     // create the kernel program
 301     const char kernel_source[] =
 302         "#include \"square.h\"\n"
 303         "#include \"div2.h\"\n"
 304         "__kernel void squareby2_kernel(__global int *x)"
 305         "{"
 306         "    x[0] = div2(square(x[0]));"
 307         "}";
 308
 309     compute::program square_program =
 310         compute::program::create_with_source(kernel_source, context);
 311
 312     std::vector<std::pair<std::string, compute::program> > header_programs;
 313     header_programs.push_back(std::make_pair("square.h", square_header_program));
 314     header_programs.push_back(std::make_pair("div2.h", div2_header_program));
 315
 316     square_program.compile("-DT=int", header_programs);
 317
 318     // link program
 319     std::vector<compute::program> programs;
 320     programs.push_back(square_program);
 321
 322     compute::program linked_program =
 323         compute::program::link(programs, context);
 324
 325     // create the square kernel
 326     compute::kernel square_kernel =
 327         linked_program.create_kernel("squareby2_kernel");
 328     BOOST_CHECK_EQUAL(square_kernel.name(), "squareby2_kernel");
 329 }
 330 #endif // BOOST_COMPUTE_CL_VERSION_1_2
 331
 332 BOOST_AUTO_TEST_CASE(build_log)
 333 {
 334     const char invalid_source[] =
 335         "__kernel void foo(__global int *input) { !@#$%^&*() }";
 336
 337     compute::program invalid_program =
 338         compute::program::create_with_source(invalid_source, context);
 339
 340     try {
 341         invalid_program.build();
 342
 343         // should not get here
 344         BOOST_CHECK(false);
 345     }
 346     catch(compute::opencl_error&){
 347         std::string log = invalid_program.build_log();
 348         BOOST_CHECK(!log.empty());
 349     }
 350 }
 351
 352 BOOST_AUTO_TEST_CASE(program_build_exception)
 353 {
 354     const char invalid_source[] =
 355         "__kernel void foo(__global int *input) { !@#$%^&*() }";
 356
 357     compute::program invalid_program =
 358         compute::program::create_with_source(invalid_source, context);
 359
 360     BOOST_CHECK_THROW(invalid_program.build(),
 361                       compute::program_build_failure);
 362
 363     try {
 364         // POCL bug: https://github.com/pocl/pocl/issues/577
 365         if(pocl_bug_issue_577(device))
 366         {
 367             invalid_program =
 368                 compute::program::create_with_source(invalid_source, context);
 369         }
 370         invalid_program.build();
 371
 372         // should not get here
 373         BOOST_CHECK(false);
 374     }
 375     catch(compute::program_build_failure& e){
 376         BOOST_CHECK(e.build_log() == invalid_program.build_log());
 377     }
 378     catch(...)
 379     {
 380         // should not get here
 381         BOOST_CHECK(false);
 382     }
 383 }
 384
 385 BOOST_AUTO_TEST_CASE(build_with_source_exception)
 386 {
 387     const char invalid_source[] =
 388         "__kernel void foo(__global int *input) { !@#$%^&*() }";
 389
 390     BOOST_CHECK_THROW(compute::program::build_with_source(invalid_source, context),
 391         compute::program_build_failure);
 392 }
 393
 394 BOOST_AUTO_TEST_CASE(build_with_source_file_exception)
 395 {
 396     std::string file_path(BOOST_COMPUTE_TEST_DATA_PATH "/invalid_program.cl");
 397     BOOST_CHECK_THROW(compute::program::build_with_source_file(file_path, context),
 398         compute::program_build_failure);
 399 }
 400
 401 BOOST_AUTO_TEST_SUITE_END()