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 #include <iostream>
12
13 #include <boost/compute/system.hpp>
14 #include <boost/compute/algorithm/copy.hpp>
15 #include <boost/compute/algorithm/accumulate.hpp>
16 #include <boost/compute/algorithm/exclusive_scan.hpp>
17 #include <boost/compute/container/vector.hpp>
18 #include <boost/compute/utility/dim.hpp>
19 #include <boost/compute/utility/source.hpp>
20
21 namespace compute = boost::compute;
22
23 const char fizz_buzz_source[] = BOOST_COMPUTE_STRINGIZE_SOURCE(
24 // returns the length of the string for the number 'n'. This is used
25 // during the first pass when we calculate the amount of space needed
26 // for each string in the fizz-buzz sequence.
27 inline uint fizz_buzz_string_length(uint n)
28 {
29 if((n % 5 == 0) && (n % 3 == 0)){
30 return sizeof("fizzbuzz");
31 }
32 else if(n % 5 == 0){
33 return sizeof("fizz");
34 }
35 else if(n % 3 == 0){
36 return sizeof("buzz");
37 }
38 else {
39 uint digits = 0;
40 while(n){
41 n /= 10;
42 digits++;
43 }
44 return digits + 1;
45 }
46 }
47
48 // first-pass kernel which calculates the string length for each number
49 // and writes it to the string_lengths array. these will then be passed
50 // to exclusive_scan() to calculate the output offsets for each string.
51 __kernel void fizz_buzz_allocate_strings(__global uint *string_lengths)
52 {
53 const uint i = get_global_id(0);
54 const uint n = i + 1;
55
56 string_lengths[i] = fizz_buzz_string_length(n);
57 }
58
59 // copy the string 's' with length 'n' to 'result' (just like strncpy())
60 inline void copy_string(__constant const char *s, uint n, __global char *result)
61 {
62 while(n--){
63 result[n] = s[n];
64 }
65 }
66
67 // reverse the string [start, end).
68 inline void reverse_string(__global char *start, __global char *end)
69 {
70 while(start < end){
71 char tmp = *end;
72 *end = *start;
73 *start = tmp;
74 start++;
75 end--;
76 }
77 }
78
79 // second-pass kernel which copies the fizz-buzz string for each number to
80 // buffer using the previously calculated offsets.
81 __kernel void fizz_buzz_copy_strings(__global const uint *offsets, __global char *buffer)
82 {
83 const uint i = get_global_id(0);
84 const uint n = i + 1;
85 const uint offset = offsets[i];
86
87 if((n % 5 == 0) && (n % 3 == 0)){
88 copy_string("fizzbuzz\n", 9, buffer + offset);
89 }
90 else if(n % 5 == 0){
91 copy_string("fizz\n", 5, buffer + offset);
92 }
93 else if(n % 3 == 0){
94 copy_string("buzz\n", 5, buffer + offset);
95 }
96 else {
97 // convert number to string and write it to the output
98 __global char *number = buffer + offset;
99 uint n_ = n;
100 while(n_){
101 *number++ = (n_%10) + '0';
102 n_ /= 10;
103 }
104 reverse_string(buffer + offset, number - 1);
105 *number = '\n';
106 }
107 }
108 );
109
main()110 int main()
111 {
112 using compute::dim;
113 using compute::uint_;
114
115 // fizz-buzz up to 100
116 size_t n = 100;
117
118 // get the default device
119 compute::device device = compute::system::default_device();
120 compute::context ctx(device);
121 compute::command_queue queue(ctx, device);
122
123 // compile the fizz-buzz program
124 compute::program fizz_buzz_program =
125 compute::program::create_with_source(fizz_buzz_source, ctx);
126 fizz_buzz_program.build();
127
128 // create a vector for the output string and computing offsets
129 compute::vector<char> output(ctx);
130 compute::vector<uint_> offsets(n, ctx);
131
132 // run the allocate kernel to calculate string lengths
133 compute::kernel allocate_kernel(fizz_buzz_program, "fizz_buzz_allocate_strings");
134 allocate_kernel.set_arg(0, offsets);
135 queue.enqueue_nd_range_kernel(allocate_kernel, dim(0), dim(n), dim(1));
136
137 // allocate space for the output string
138 output.resize(
139 compute::accumulate(offsets.begin(), offsets.end(), 0, queue)
140 );
141
142 // scan string lengths for each number to calculate the output offsets
143 compute::exclusive_scan(
144 offsets.begin(), offsets.end(), offsets.begin(), queue
145 );
146
147 // run the copy kernel to fill the output buffer
148 compute::kernel copy_kernel(fizz_buzz_program, "fizz_buzz_copy_strings");
149 copy_kernel.set_arg(0, offsets);
150 copy_kernel.set_arg(1, output);
151 queue.enqueue_nd_range_kernel(copy_kernel, dim(0), dim(n), dim(1));
152
153 // copy the string to the host and print it to stdout
154 std::string str;
155 str.resize(output.size());
156 compute::copy(output.begin(), output.end(), str.begin(), queue);
157 std::cout << str;
158
159 return 0;
160 }
161