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 TestLocalBuffer
12 #include <boost/test/unit_test.hpp>
13
14 #include <iostream>
15
16 #include <boost/compute/core.hpp>
17 #include <boost/compute/memory/local_buffer.hpp>
18 #include <boost/compute/utility/source.hpp>
19
20 #include "context_setup.hpp"
21
22 namespace compute = boost::compute;
23
BOOST_AUTO_TEST_CASE(local_buffer_arg)24 BOOST_AUTO_TEST_CASE(local_buffer_arg)
25 {
26 if(device.get_info<CL_DEVICE_LOCAL_MEM_TYPE>() != CL_LOCAL){
27 std::cerr << "skipping local buffer arg test: "
28 << "device does not support local memory" << std::endl;
29 return;
30 }
31
32 const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
33 __kernel void foo(__local float *local_buffer,
34 __global float *global_buffer)
35 {
36 const uint gid = get_global_id(0);
37 const uint lid = get_local_id(0);
38
39 local_buffer[lid] = gid;
40
41 global_buffer[gid] = local_buffer[lid];
42 }
43 );
44
45 // create and build program
46 compute::program program =
47 compute::program::build_with_source(source, context);
48
49 // create kernel
50 compute::kernel kernel = program.create_kernel("foo");
51
52 // setup kernel arguments
53 compute::buffer global_buf(context, 128 * sizeof(float));
54 kernel.set_arg(0, compute::local_buffer<float>(32));
55 kernel.set_arg(1, global_buf);
56
57 // some implementations don't correctly report dynamically-set local buffer sizes
58 if(kernel.get_work_group_info<cl_ulong>(device, CL_KERNEL_LOCAL_MEM_SIZE) == 0){
59 std::cerr << "skipping checks for local memory size, device reports "
60 << "zero after setting dynamically-sized local buffer size"
61 << std::endl;
62 return;
63 }
64
65 // check actual memory size
66 BOOST_CHECK_GE(
67 kernel.get_work_group_info<cl_ulong>(device, CL_KERNEL_LOCAL_MEM_SIZE),
68 32 * sizeof(float)
69 );
70
71 // increase local buffer size and check new actual local memory size
72 kernel.set_arg(0, compute::local_buffer<float>(64));
73
74 BOOST_CHECK_GE(
75 kernel.get_work_group_info<cl_ulong>(device, CL_KERNEL_LOCAL_MEM_SIZE),
76 64 * sizeof(float)
77 );
78 }
79
80 BOOST_AUTO_TEST_SUITE_END()
81