Skip to content

Commit 11401ed

Browse files
committed
Allows passing struct by value, utilizing constant memory on the device
1 parent 36c8913 commit 11401ed

File tree

2 files changed

+41
-0
lines changed

2 files changed

+41
-0
lines changed

include/boost/compute/types/struct.hpp

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -168,6 +168,14 @@ inline std::string adapt_struct_insert_member(T (Struct::*)[N], const char *name
168168
) \
169169
<< "}"; \
170170
} \
171+
template<> \
172+
struct set_kernel_arg<type> \
173+
{ \
174+
void operator()(kernel &kernel_, size_t index, const type &c) \
175+
{ \
176+
kernel_.set_arg(index, sizeof(type), &c); \
177+
} \
178+
}; \
171179
}}}
172180

173181
#endif // BOOST_COMPUTE_TYPES_STRUCT_HPP

test/test_struct.cpp

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -23,6 +23,8 @@
2323
#include <boost/compute/type_traits/type_definition.hpp>
2424
#include <boost/compute/utility/source.hpp>
2525

26+
#include "check_macros.hpp"
27+
2628
namespace compute = boost::compute;
2729

2830
// example code defining an atom class
@@ -131,6 +133,37 @@ BOOST_AUTO_TEST_CASE(custom_kernel)
131133
queue.enqueue_1d_range_kernel(custom_kernel, 0, atoms.size(), 1);
132134
}
133135

136+
BOOST_AUTO_TEST_CASE(custom_kernel_set_struct_by_value)
137+
{
138+
std::string source = BOOST_COMPUTE_STRINGIZE_SOURCE(
139+
__kernel void custom_kernel(Atom atom,
140+
__global float *position,
141+
__global int *number)
142+
{
143+
position[0] = atom.x;
144+
position[1] = atom.y;
145+
position[2] = atom.z;
146+
number[0] = atom.number;
147+
}
148+
);
149+
source = compute::type_definition<chemistry::Atom>() + "\n" + source;
150+
compute::program program =
151+
compute::program::build_with_source(source, context);
152+
compute::kernel custom_kernel = program.create_kernel("custom_kernel");
153+
154+
chemistry::Atom atom(1.0f, 2.0f, 3.0f, 4);
155+
compute::vector<float> position(3);
156+
compute::vector<int> number(1);
157+
158+
custom_kernel.set_arg(0, atom);
159+
custom_kernel.set_arg(1, position);
160+
custom_kernel.set_arg(2, number);
161+
queue.enqueue_task(custom_kernel);
162+
163+
CHECK_RANGE_EQUAL(float, 3, position, (1.0f, 2.0f, 3.0f));
164+
BOOST_CHECK_EQUAL(number.at(0), 4);
165+
}
166+
134167
// Creates a StructWithArray containing 'x', 'y', 'z'.
135168
StructWithArray make_struct_with_array(int x, int y, int z)
136169
{

0 commit comments

Comments
 (0)