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