xref: /aosp_15_r20/external/OpenCL-CLHPP/examples/src/trivialSizeTCompat.cpp (revision 6fee86a4f833e4f32f25770a262884407554133d)
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