1 #define CL_HPP_ENABLE_EXCEPTIONS
2 #define CL_HPP_TARGET_OPENCL_VERSION 200
3 #define CL_HPP_ENABLE_SIZE_T_COMPATIBILITY
4
5 #include <CL/opencl.hpp>
6 #include <iostream>
7 #include <vector>
8
9 const int numElements = 32;
10
main(void)11 int main(void)
12 {
13 // Filter for a 2.0 platform and set it as the default
14 std::vector<cl::Platform> platforms;
15 cl::Platform::get(&platforms);
16 cl::Platform plat;
17 for (auto &p : platforms) {
18 std::string platver = p.getInfo<CL_PLATFORM_VERSION>();
19 if (platver.find("OpenCL 2.") != std::string::npos) {
20 plat = p;
21 }
22 }
23 if (plat() == 0) {
24 std::cout << "No OpenCL 2.0 platform found.\n";
25 return -1;
26 }
27
28 cl::Platform newP = cl::Platform::setDefault(plat);
29 if (newP != plat) {
30 std::cout << "Error setting default platform.";
31 return -1;
32 }
33 cl::Program vectorAddProgram(
34 std::string(
35 "global int globalA;"
36 "kernel void updateGlobal(){"
37 " globalA = 75;"
38 "}"
39 "kernel void vectorAdd(global const int *inputA, global const int *inputB, global int *output, int val, write_only pipe int outPipe){"
40 " output[get_global_id(0)] = inputA[get_global_id(0)] + inputB[get_global_id(0)] + val;"
41 " write_pipe(outPipe, &val);"
42 " queue_t default_queue = get_default_queue(); "
43 " ndrange_t ndrange = ndrange_1D(get_global_size(0), get_global_size(0)); "
44 " enqueue_kernel(default_queue, CLK_ENQUEUE_FLAGS_WAIT_KERNEL, ndrange, "
45 " ^{"
46 " output[get_global_size(0)+get_global_id(0)] = inputA[get_global_size(0)+get_global_id(0)] + inputB[get_global_size(0)+get_global_id(0)] + globalA;"
47 " });"
48 "}")
49 , false);
50 try {
51 vectorAddProgram.build("-cl-std=CL2.0");
52 }
53 catch (...) {
54 std::string bl = vectorAddProgram.getBuildInfo<CL_PROGRAM_BUILD_LOG>(cl::Device::getDefault());
55 std::cerr << bl << std::endl;
56 }
57
58 // Get and run kernel that initializes the program-scope global
59 // A test for kernels that take no arguments
60 auto program2Kernel =
61 cl::KernelFunctor<>(vectorAddProgram, "updateGlobal");
62 program2Kernel(
63 cl::EnqueueArgs(
64 cl::NDRange(1)));
65
66 auto vectorAddKernel =
67 cl::KernelFunctor<
68 cl::Buffer&,
69 cl::Buffer&,
70 cl::Buffer&,
71 int,
72 cl::Pipe&
73 >(vectorAddProgram, "vectorAdd");
74
75 std::vector<int> inputA(numElements, 1);
76 std::vector<int> inputB(numElements, 2);
77 std::vector<int> output(numElements, 0xdeadbeef);
78 cl::Buffer inputABuffer(inputA.begin(), inputA.end(), true);
79 cl::Buffer inputBBuffer(inputB.begin(), inputB.end(), true);
80 cl::Buffer outputBuffer(output.begin(), output.end(), false);
81 cl::Pipe aPipe(sizeof(cl_int), numElements / 2);
82 // Unfortunately, there is no way to check for a default or know if a kernel needs one
83 // so the user has to create one
84 // We can't preemptively do so on device creation because they cannot then replace it
85 cl::DeviceCommandQueue deviceQueue = cl::DeviceCommandQueue::makeDefault(
86 cl::Context::getDefault(), cl::Device::getDefault());
87
88 vectorAddKernel(
89 cl::EnqueueArgs(
90 cl::NDRange(numElements/2),
91 cl::NDRange(numElements/2)),
92 inputABuffer,
93 inputBBuffer,
94 outputBuffer,
95 3,
96 aPipe);
97
98 cl_int error;
99 vectorAddKernel(
100 cl::EnqueueArgs(
101 cl::NDRange(numElements/2),
102 cl::NDRange(numElements/2)),
103 inputABuffer,
104 inputBBuffer,
105 outputBuffer,
106 3,
107 aPipe,
108 error);
109
110 cl::array<size_t, 3> WGSizeResultArray = vectorAddKernel.getKernel().getWorkGroupInfo<CL_KERNEL_COMPILE_WORK_GROUP_SIZE>(cl::Device::getDefault());
111 std::cout << "Array return: " << WGSizeResultArray[0] << ", " << WGSizeResultArray[1] << ", " << WGSizeResultArray[2] << "\n";
112 cl::size_t<3> WGSizeResult = vectorAddKernel.getKernel().getWorkGroupInfo<CL_KERNEL_COMPILE_WORK_GROUP_SIZE>(cl::Device::getDefault());
113 std::cout << "Size_t return: " << WGSizeResult[0] << ", " << WGSizeResult[1] << ", " << WGSizeResult[2] << "\n";
114
115 cl::copy(outputBuffer, output.begin(), output.end());
116
117 cl::Device d = cl::Device::getDefault();
118 std::cout << "Max pipe args: " << d.getInfo<CL_DEVICE_MAX_PIPE_ARGS>() << "\n";
119 std::cout << "Max pipe active reservations: " << d.getInfo<CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS>() << "\n";
120 std::cout << "Max pipe packet size: " << d.getInfo<CL_DEVICE_PIPE_MAX_PACKET_SIZE>() << "\n";
121
122
123
124 std::cout << "Output:\n";
125 for (int i = 1; i < numElements; ++i) {
126 std::cout << "\t" << output[i] << "\n";
127 }
128 std::cout << "\n\n";
129
130 }
131