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 TestAmdCppKernel
12 #include <boost/test/unit_test.hpp>
13
14 #include <iostream>
15
16 #include <boost/compute/kernel.hpp>
17 #include <boost/compute/program.hpp>
18 #include <boost/compute/system.hpp>
19 #include <boost/compute/container/vector.hpp>
20 #include <boost/compute/detail/vendor.hpp>
21 #include <boost/compute/utility/source.hpp>
22
23 #include "check_macros.hpp"
24 #include "context_setup.hpp"
25
26 namespace compute = boost::compute;
27
BOOST_AUTO_TEST_CASE(amd_template_function)28 BOOST_AUTO_TEST_CASE(amd_template_function)
29 {
30 if(!compute::detail::is_amd_device(device)){
31 std::cerr << "skipping amd_template_function test: c++ static kernel "
32 "language is only supported on AMD devices." << std::endl;
33 return;
34 }
35
36 const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
37 template<typename T>
38 inline T square(const T x)
39 {
40 return x * x;
41 }
42
43 template<typename T>
44 __kernel void square_kernel(__global T *data)
45 {
46 const uint i = get_global_id(0);
47 data[i] = square(data[i]);
48 }
49
50 template __attribute__((mangled_name(square_kernel_int)))
51 __kernel void square_kernel(__global int *data);
52 );
53
54 int data[] = { 1, 2, 3, 4 };
55 compute::vector<int> vec(data, data + 4, queue);
56
57 compute::program square_program =
58 compute::program::build_with_source(source, context, "-x clc++");
59
60 compute::kernel square_kernel(square_program, "square_kernel_int");
61 square_kernel.set_arg(0, vec);
62
63 queue.enqueue_1d_range_kernel(square_kernel, 0, vec.size(), 4);
64
65 CHECK_RANGE_EQUAL(int, 4, vec, (1, 4, 9, 16));
66 }
67
68 BOOST_AUTO_TEST_SUITE_END()
69