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

   1 //---------------------------------------------------------------------------//
   2 // Copyright (c) 2013-2014 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 TestSvmPtr
  12 #include <boost/test/unit_test.hpp>
  13
  14 #include <iostream>
  15
  16 #include <boost/compute/core.hpp>
  17 #include <boost/compute/svm.hpp>
  18 #include <boost/compute/container/vector.hpp>
  19 #include <boost/compute/utility/source.hpp>
  20
  21 #include "quirks.hpp"
  22 #include "check_macros.hpp"
  23 #include "context_setup.hpp"
  24
  25 namespace compute = boost::compute;
  26
  27 BOOST_AUTO_TEST_CASE(empty)
  28 {
  29 }
  30
  31 #ifdef BOOST_COMPUTE_CL_VERSION_2_0
  32 BOOST_AUTO_TEST_CASE(alloc)
  33 {
  34     REQUIRES_OPENCL_VERSION(2, 0);
  35
  36     compute::svm_ptr<cl_int> ptr = compute::svm_alloc<cl_int>(context, 8);
  37     compute::svm_free(context, ptr);
  38 }
  39
  40 BOOST_AUTO_TEST_CASE(svmmemcpy)
  41 {
  42     REQUIRES_OPENCL_VERSION(2, 0);
  43
  44     if(bug_in_svmmemcpy(device)){
  45         std::cerr << "skipping svmmemcpy test case" << std::endl;
  46         return;
  47     }
  48
  49     cl_int input[] = { 1, 2, 3, 4, 5, 6, 7, 8 };
  50     cl_int output[] = { 0, 0, 0, 0, 0, 0, 0, 0 };
  51     compute::svm_ptr<cl_int> ptr = compute::svm_alloc<cl_int>(context, 8);
  52     compute::svm_ptr<cl_int> ptr2 = compute::svm_alloc<cl_int>(context, 8);
  53
  54     // copying from and to host mem
  55     queue.enqueue_svm_memcpy(ptr.get(), input, 8 * sizeof(cl_int));
  56     queue.enqueue_svm_memcpy(output, ptr.get(), 8 * sizeof(cl_int));
  57     queue.finish();
  58
  59     CHECK_HOST_RANGE_EQUAL(cl_int, 8, output, (1, 2, 3, 4, 5, 6, 7, 8));
  60
  61     // copying between svm mem
  62     queue.enqueue_svm_memcpy(ptr2.get(), ptr.get(), 8 * sizeof(cl_int));
  63     queue.enqueue_svm_memcpy(output, ptr2.get(), 8 * sizeof(cl_int));
  64     queue.finish();
  65
  66     CHECK_HOST_RANGE_EQUAL(cl_int, 8, output, (1, 2, 3, 4, 5, 6, 7, 8));
  67
  68     compute::svm_free(context, ptr);
  69     compute::svm_free(context, ptr2);
  70 }
  71
  72 BOOST_AUTO_TEST_CASE(sum_svm_kernel)
  73 {
  74     REQUIRES_OPENCL_VERSION(2, 0);
  75
  76     const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
  77         __kernel void sum_svm_mem(__global const int *ptr, __global int *result)
  78         {
  79             int sum = 0;
  80             for(uint i = 0; i < 8; i++){
  81                 sum += ptr[i];
  82             }
  83             *result = sum;
  84         }
  85     );
  86
  87     compute::program program =
  88         compute::program::build_with_source(source, context, "-cl-std=CL2.0");
  89
  90     compute::kernel sum_svm_mem_kernel = program.create_kernel("sum_svm_mem");
  91
  92     cl_int data[] = { 1, 2, 3, 4, 5, 6, 7, 8 };
  93     compute::svm_ptr<cl_int> ptr = compute::svm_alloc<cl_int>(context, 8);
  94     queue.enqueue_svm_map(ptr.get(), 8 * sizeof(cl_int), CL_MAP_WRITE);
  95     for(size_t i = 0; i < 8; i ++) {
  96         static_cast<cl_int*>(ptr.get())[i] = data[i];
  97     }
  98     queue.enqueue_svm_unmap(ptr.get());
  99
 100     compute::vector<cl_int> result(1, context);
 101
 102     sum_svm_mem_kernel.set_arg(0, ptr);
 103     sum_svm_mem_kernel.set_arg(1, result);
 104     queue.enqueue_task(sum_svm_mem_kernel);
 105
 106     queue.finish();
 107     BOOST_CHECK_EQUAL(result[0], (36));
 108
 109     compute::svm_free(context, ptr);
 110 }
 111 #endif // BOOST_COMPUTE_CL_VERSION_2_0
 112
 113 #ifdef BOOST_COMPUTE_CL_VERSION_2_1
 114 BOOST_AUTO_TEST_CASE(migrate)
 115 {
 116     REQUIRES_OPENCL_VERSION(2, 1);
 117
 118     compute::svm_ptr<cl_int> ptr =
 119         compute::svm_alloc<cl_int>(context, 8);
 120
 121     // Migrate to device
 122     std::vector<const void*> ptrs(1, ptr.get());
 123     std::vector<size_t> sizes(1, 8 * sizeof(cl_int));
 124     queue.enqueue_svm_migrate_memory(ptrs, sizes).wait();
 125
 126     // Set on device
 127     const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
 128         __kernel void foo(__global int *ptr)
 129         {
 130             for(int i = 0; i < 8; i++){
 131                 ptr[i] = i;
 132             }
 133         }
 134     );
 135     compute::program program =
 136         compute::program::build_with_source(source, context, "-cl-std=CL2.0");
 137     compute::kernel foo_kernel = program.create_kernel("foo");
 138     foo_kernel.set_arg(0, ptr);
 139     queue.enqueue_task(foo_kernel).wait();
 140
 141     // Migrate to host
 142     queue.enqueue_svm_migrate_memory(
 143         ptr.get(), 0, boost::compute::command_queue::migrate_to_host
 144     ).wait();
 145
 146     // Check
 147     CHECK_HOST_RANGE_EQUAL(
 148         cl_int, 8,
 149         static_cast<cl_int*>(ptr.get()),
 150         (0, 1, 2, 3, 4, 5, 6, 7)
 151     );
 152     compute::svm_free(context, ptr);
 153 }
 154 #endif // BOOST_COMPUTE_CL_VERSION_2_1
 155
 156 BOOST_AUTO_TEST_SUITE_END()