1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
|
//---------------------------------------------------------------------------//
// Copyright (c) 2013-2014 Kyle Lutz <kyle.r.lutz@gmail.com>
//
// Distributed under the Boost Software License, Version 1.0
// See accompanying file LICENSE_1_0.txt or copy at
// http://www.boost.org/LICENSE_1_0.txt
//
// See http://boostorg.github.com/compute for more information.
//---------------------------------------------------------------------------//
#define BOOST_TEST_MODULE TestLocalBuffer
#include <boost/test/unit_test.hpp>
#include <iostream>
#include <boost/compute/core.hpp>
#include <boost/compute/memory/local_buffer.hpp>
#include <boost/compute/utility/source.hpp>
#include "context_setup.hpp"
namespace compute = boost::compute;
BOOST_AUTO_TEST_CASE(local_buffer_arg)
{
if(device.get_info<CL_DEVICE_LOCAL_MEM_TYPE>() != CL_LOCAL){
std::cerr << "skipping local buffer arg test: "
<< "device does not support local memory" << std::endl;
return;
}
const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
__kernel void foo(__local float *local_buffer,
__global float *global_buffer)
{
const uint gid = get_global_id(0);
const uint lid = get_local_id(0);
local_buffer[lid] = gid;
global_buffer[gid] = local_buffer[lid];
}
);
// create and build program
compute::program program =
compute::program::build_with_source(source, context);
// create kernel
compute::kernel kernel = program.create_kernel("foo");
// setup kernel arguments
compute::buffer global_buf(context, 128 * sizeof(float));
kernel.set_arg(0, compute::local_buffer<float>(32));
kernel.set_arg(1, global_buf);
// some implementations don't correctly report dynamically-set local buffer sizes
if(kernel.get_work_group_info<cl_ulong>(device, CL_KERNEL_LOCAL_MEM_SIZE) == 0){
std::cerr << "skipping checks for local memory size, device reports "
<< "zero after setting dynamically-sized local buffer size"
<< std::endl;
return;
}
// check actual memory size
BOOST_CHECK_GE(
kernel.get_work_group_info<cl_ulong>(device, CL_KERNEL_LOCAL_MEM_SIZE),
32 * sizeof(float)
);
// increase local buffer size and check new actual local memory size
kernel.set_arg(0, compute::local_buffer<float>(64));
BOOST_CHECK_GE(
kernel.get_work_group_info<cl_ulong>(device, CL_KERNEL_LOCAL_MEM_SIZE),
64 * sizeof(float)
);
}
BOOST_AUTO_TEST_SUITE_END()
|
