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