|
| 1 | +/* |
| 2 | +// Copyright (c) 2019-2026 Ben Ashbaugh |
| 3 | +// |
| 4 | +// SPDX-License-Identifier: MIT |
| 5 | +*/ |
| 6 | + |
| 7 | +#include <popl/popl.hpp> |
| 8 | + |
| 9 | +#include <CL/opencl.hpp> |
| 10 | + |
| 11 | +#include "util.hpp" |
| 12 | + |
| 13 | +static const char kernelString[] = R"CLC( |
| 14 | +kernel void test(global uint* dst) |
| 15 | +{ |
| 16 | + uint id = get_global_id(0); |
| 17 | + dst[id] = id; |
| 18 | +} |
| 19 | +)CLC"; |
| 20 | + |
| 21 | +#if !defined(cl_khr_suggested_local_work_size) |
| 22 | +#pragma message("suggestedlocalworksize: cl_khr_suggested_local_work_size. Please update your OpenCL headers.") |
| 23 | +#endif |
| 24 | + |
| 25 | +int main(int argc, char** argv) |
| 26 | +{ |
| 27 | + int platformIndex = 0; |
| 28 | + int deviceIndex = 0; |
| 29 | + |
| 30 | + size_t gws = 1024*1024; |
| 31 | + |
| 32 | + { |
| 33 | + popl::OptionParser op("Supported Options"); |
| 34 | + op.add<popl::Value<int>>("p", "platform", "Platform Index", platformIndex, &platformIndex); |
| 35 | + op.add<popl::Value<int>>("d", "device", "Device Index", deviceIndex, &deviceIndex); |
| 36 | + op.add<popl::Value<size_t>>("g", "gws", "Global Work Size", gws, &gws); |
| 37 | + |
| 38 | + bool printUsage = false; |
| 39 | + try { |
| 40 | + op.parse(argc, argv); |
| 41 | + } catch (std::exception& e) { |
| 42 | + fprintf(stderr, "Error: %s\n\n", e.what()); |
| 43 | + printUsage = true; |
| 44 | + } |
| 45 | + |
| 46 | + if (printUsage || !op.unknown_options().empty() || !op.non_option_args().empty()) { |
| 47 | + fprintf(stderr, |
| 48 | + "Usage: suggestedlocalworksize [options]\n" |
| 49 | + "%s", op.help().c_str()); |
| 50 | + return -1; |
| 51 | + } |
| 52 | + } |
| 53 | + |
| 54 | + std::vector<cl::Platform> platforms; |
| 55 | + cl::Platform::get(&platforms); |
| 56 | + |
| 57 | + if (!checkPlatformIndex(platforms, platformIndex)) { |
| 58 | + return -1; |
| 59 | + } |
| 60 | + |
| 61 | + printf("Running on platform: %s\n", |
| 62 | + platforms[platformIndex].getInfo<CL_PLATFORM_NAME>().c_str() ); |
| 63 | + |
| 64 | + std::vector<cl::Device> devices; |
| 65 | + platforms[platformIndex].getDevices(CL_DEVICE_TYPE_ALL, &devices); |
| 66 | + |
| 67 | + printf("Running on device: %s\n", |
| 68 | + devices[deviceIndex].getInfo<CL_DEVICE_NAME>().c_str() ); |
| 69 | + |
| 70 | + if (checkDeviceForExtension(devices[deviceIndex], CL_KHR_SUGGESTED_LOCAL_WORK_SIZE_EXTENSION_NAME)) { |
| 71 | + printf("Device supports %s.\n", CL_KHR_SUGGESTED_LOCAL_WORK_SIZE_EXTENSION_NAME); |
| 72 | + } else { |
| 73 | + printf("Device does not support %s, exiting.\n", CL_KHR_SUGGESTED_LOCAL_WORK_SIZE_EXTENSION_NAME); |
| 74 | + return -1; |
| 75 | + } |
| 76 | + |
| 77 | + clGetKernelSuggestedLocalWorkSizeKHR_fn clGetKernelSuggestedLocalWorkSizeKHR = |
| 78 | + (clGetKernelSuggestedLocalWorkSizeKHR_fn)clGetExtensionFunctionAddressForPlatform( |
| 79 | + platforms[platformIndex](), |
| 80 | + "clGetKernelSuggestedLocalWorkSizeKHR" ); |
| 81 | + if (clGetKernelSuggestedLocalWorkSizeKHR == nullptr) { |
| 82 | + fprintf(stderr, "Failed to get function pointer for clGetKernelSuggestedLocalWorkSizeKHR, exiting.\n"); |
| 83 | + return -1; |
| 84 | + } |
| 85 | + |
| 86 | + cl::Context context{devices[deviceIndex]}; |
| 87 | + cl::CommandQueue commandQueue{context, devices[deviceIndex]}; |
| 88 | + |
| 89 | + cl::Program program{ context, kernelString }; |
| 90 | + program.build(); |
| 91 | + cl::Kernel kernel = cl::Kernel{ program, "test" }; |
| 92 | + |
| 93 | + cl::Buffer buf = cl::Buffer{ |
| 94 | + context, |
| 95 | + CL_MEM_ALLOC_HOST_PTR, |
| 96 | + gws * sizeof( cl_uint ) }; |
| 97 | + |
| 98 | + kernel.setArg(0, buf); |
| 99 | + |
| 100 | + size_t suggestedLocalWorkSize = 0; |
| 101 | + cl_int errorCode = clGetKernelSuggestedLocalWorkSizeKHR( |
| 102 | + commandQueue(), |
| 103 | + kernel(), |
| 104 | + 1, |
| 105 | + nullptr, |
| 106 | + &gws, |
| 107 | + &suggestedLocalWorkSize ); |
| 108 | + |
| 109 | + if (errorCode != CL_SUCCESS) { |
| 110 | + fprintf(stderr, "clGetKernelSuggestedLocalWorkSizeKHR failed with error code %d, exiting.\n", errorCode); |
| 111 | + return -1; |
| 112 | + } |
| 113 | + |
| 114 | + printf("Suggested local work size for global work size %zu is: %zu\n", gws, suggestedLocalWorkSize); |
| 115 | + return 0; |
| 116 | +} |
0 commit comments