File: test_amd_cpp_kernel_language.cpp

package info (click to toggle)
boost1.83 1.83.0-5
  • links: PTS, VCS
  • area: main
  • in suites: forky, sid
  • size: 545,632 kB
  • sloc: cpp: 3,857,086; xml: 125,552; ansic: 34,414; python: 25,887; asm: 5,276; sh: 4,799; ada: 1,681; makefile: 1,629; perl: 1,212; pascal: 1,139; sql: 810; yacc: 478; ruby: 102; lisp: 24; csh: 6
file content (68 lines) | stat: -rw-r--r-- 2,108 bytes parent folder | download | duplicates (16)
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
//---------------------------------------------------------------------------//
// 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 TestAmdCppKernel
#include <boost/test/unit_test.hpp>

#include <iostream>

#include <boost/compute/kernel.hpp>
#include <boost/compute/program.hpp>
#include <boost/compute/system.hpp>
#include <boost/compute/container/vector.hpp>
#include <boost/compute/detail/vendor.hpp>
#include <boost/compute/utility/source.hpp>

#include "check_macros.hpp"
#include "context_setup.hpp"

namespace compute = boost::compute;

BOOST_AUTO_TEST_CASE(amd_template_function)
{
    if(!compute::detail::is_amd_device(device)){
        std::cerr << "skipping amd_template_function test: c++ static kernel "
                     "language is only supported on AMD devices." << std::endl;
        return;
    }

    const char source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
        template<typename T>
        inline T square(const T x)
        {
            return x * x;
        }

        template<typename T>
        __kernel void square_kernel(__global T *data)
        {
            const uint i = get_global_id(0);
            data[i] = square(data[i]);
        }

        template __attribute__((mangled_name(square_kernel_int)))
        __kernel void square_kernel(__global int *data);
    );

    int data[] = { 1, 2, 3, 4 };
    compute::vector<int> vec(data, data + 4, queue);

    compute::program square_program =
        compute::program::build_with_source(source, context, "-x clc++");

    compute::kernel square_kernel(square_program, "square_kernel_int");
    square_kernel.set_arg(0, vec);

    queue.enqueue_1d_range_kernel(square_kernel, 0, vec.size(), 4);

    CHECK_RANGE_EQUAL(int, 4, vec, (1, 4, 9, 16));
}

BOOST_AUTO_TEST_SUITE_END()