1 // 2 // Copyright (c) 2017 The Khronos Group Inc. 3 // 4 // Licensed under the Apache License, Version 2.0 (the "License"); 5 // you may not use this file except in compliance with the License. 6 // You may obtain a copy of the License at 7 // 8 // http://www.apache.org/licenses/LICENSE-2.0 9 // 10 // Unless required by applicable law or agreed to in writing, software 11 // distributed under the License is distributed on an "AS IS" BASIS, 12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. 13 // See the License for the specific language governing permissions and 14 // limitations under the License. 15 // 16 #ifndef _kernelHelpers_h 17 #define _kernelHelpers_h 18 19 // Configuration 20 #include "../config.hpp" 21 22 #include "compat.h" 23 #include "testHarness.h" 24 25 #include <stdio.h> 26 #include <stdlib.h> 27 28 #if defined(__MINGW32__) 29 #include <malloc.h> 30 #endif 31 32 #include <string.h> 33 34 #ifdef __APPLE__ 35 #include <OpenCL/opencl.h> 36 #else 37 #include <CL/opencl.h> 38 #endif 39 40 #include "deviceInfo.h" 41 #include "harness/alloc.h" 42 43 #include <functional> 44 45 /* 46 * The below code is intended to be used at the top of kernels that appear 47 * inline in files to set line and file info for the kernel: 48 * 49 * const char *source = { 50 * INIT_OPENCL_DEBUG_INFO 51 * "__kernel void foo( int x )\n" 52 * "{\n" 53 * " ...\n" 54 * "}\n" 55 * }; 56 */ 57 #define INIT_OPENCL_DEBUG_INFO SET_OPENCL_LINE_INFO(__LINE__, __FILE__) 58 #define SET_OPENCL_LINE_INFO(_line, _file) \ 59 "#line " STRINGIFY(_line) " " STRINGIFY(_file) "\n" 60 #ifndef STRINGIFY_VALUE 61 #define STRINGIFY_VALUE(_x) STRINGIFY(_x) 62 #endif 63 #ifndef STRINGIFY 64 #define STRINGIFY(_x) #_x 65 #endif 66 67 const int MAX_LEN_FOR_KERNEL_LIST = 20; 68 69 /* Helper that creates a single program and kernel from a single-kernel program 70 * source */ 71 extern int 72 create_single_kernel_helper(cl_context context, cl_program *outProgram, 73 cl_kernel *outKernel, unsigned int numKernelLines, 74 const char **kernelProgram, const char *kernelName, 75 const char *buildOptions = NULL); 76 77 extern int create_single_kernel_helper_with_build_options( 78 cl_context context, cl_program *outProgram, cl_kernel *outKernel, 79 unsigned int numKernelLines, const char **kernelProgram, 80 const char *kernelName, const char *buildOptions); 81 82 extern int create_single_kernel_helper_create_program( 83 cl_context context, cl_program *outProgram, unsigned int numKernelLines, 84 const char **kernelProgram, const char *buildOptions = NULL); 85 86 extern int create_single_kernel_helper_create_program_for_device( 87 cl_context context, cl_device_id device, cl_program *outProgram, 88 unsigned int numKernelLines, const char **kernelProgram, 89 const char *buildOptions = NULL); 90 91 /* Creates OpenCL C++ program. This one must be used for creating OpenCL C++ 92 * program. */ 93 extern int create_openclcpp_program(cl_context context, cl_program *outProgram, 94 unsigned int numKernelLines, 95 const char **kernelProgram, 96 const char *buildOptions = NULL); 97 98 /* Builds program (outProgram) and creates one kernel */ 99 int build_program_create_kernel_helper( 100 cl_context context, cl_program *outProgram, cl_kernel *outKernel, 101 unsigned int numKernelLines, const char **kernelProgram, 102 const char *kernelName, const char *buildOptions = NULL); 103 104 /* Helper to obtain the biggest fit work group size for all the devices in a 105 * given group and for the given global thread size */ 106 extern int get_max_common_work_group_size(cl_context context, cl_kernel kernel, 107 size_t globalThreadSize, 108 size_t *outSize); 109 110 /* Helper to obtain the biggest fit work group size for all the devices in a 111 * given group and for the given global thread size */ 112 extern int get_max_common_2D_work_group_size(cl_context context, 113 cl_kernel kernel, 114 size_t *globalThreadSize, 115 size_t *outSizes); 116 117 /* Helper to obtain the biggest fit work group size for all the devices in a 118 * given group and for the given global thread size */ 119 extern int get_max_common_3D_work_group_size(cl_context context, 120 cl_kernel kernel, 121 size_t *globalThreadSize, 122 size_t *outSizes); 123 124 /* Helper to obtain the biggest allowed work group size for all the devices in a 125 * given group */ 126 extern int get_max_allowed_work_group_size(cl_context context, cl_kernel kernel, 127 size_t *outSize, size_t *outLimits); 128 129 /* Helper to obtain the biggest allowed 1D work group size on a given device */ 130 extern int get_max_allowed_1d_work_group_size_on_device(cl_device_id device, 131 cl_kernel kernel, 132 size_t *outSize); 133 134 /* Helper to determine if a device supports an image format */ 135 extern int is_image_format_supported(cl_context context, cl_mem_flags flags, 136 cl_mem_object_type image_type, 137 const cl_image_format *fmt); 138 139 /* Helper to get pixel size for a pixel format */ 140 size_t get_pixel_bytes(const cl_image_format *fmt); 141 142 /* Verify the given device supports images. */ 143 extern test_status verifyImageSupport(cl_device_id device); 144 145 /* Checks that the given device supports images. Same as verify, but doesn't 146 * print an error */ 147 extern int checkForImageSupport(cl_device_id device); 148 extern int checkFor3DImageSupport(cl_device_id device); 149 extern int checkForReadWriteImageSupport(cl_device_id device); 150 151 /* Checks that a given queue property is supported on the specified device. 152 * Returns 1 if supported, 0 if not or an error. */ 153 extern int checkDeviceForQueueSupport(cl_device_id device, 154 cl_command_queue_properties prop); 155 156 /* Helper to obtain the min alignment for a given context, i.e the max of all 157 * min alignments for devices attached to the context*/ 158 size_t get_min_alignment(cl_context context); 159 160 /* Helper to obtain the default rounding mode for single precision computation. 161 * (Double is always CL_FP_ROUND_TO_NEAREST.) Returns 0 on error. */ 162 cl_device_fp_config get_default_rounding_mode(cl_device_id device); 163 164 #define PASSIVE_REQUIRE_IMAGE_SUPPORT(device) \ 165 if (checkForImageSupport(device)) \ 166 { \ 167 log_info( \ 168 "\n\tNote: device does not support images. Skipping test...\n"); \ 169 return TEST_SKIPPED_ITSELF; \ 170 } 171 172 #define PASSIVE_REQUIRE_3D_IMAGE_SUPPORT(device) \ 173 if (checkFor3DImageSupport(device)) \ 174 { \ 175 log_info("\n\tNote: device does not support 3D images. Skipping " \ 176 "test...\n"); \ 177 return TEST_SKIPPED_ITSELF; \ 178 } 179 180 #define PASSIVE_REQUIRE_FP16_SUPPORT(device) \ 181 if (!device_supports_half(device)) \ 182 { \ 183 log_info( \ 184 "\n\tNote: device does not support fp16. Skipping test...\n"); \ 185 return TEST_SKIPPED_ITSELF; \ 186 } 187 188 /* Prints out the standard device header for all tests given the device to print 189 * for */ 190 extern int printDeviceHeader(cl_device_id device); 191 192 // Execute the CL_DEVICE_OPENCL_C_VERSION query and return the OpenCL C version 193 // is supported by the device. 194 Version get_device_cl_c_version(cl_device_id device); 195 196 // Gets the latest (potentially non-backward compatible) OpenCL C version 197 // supported by the device. 198 Version get_device_latest_cl_c_version(cl_device_id device); 199 200 // Gets the maximum universally supported OpenCL C version in a context, i.e. 201 // the OpenCL C version supported by all devices in a context. 202 Version get_max_OpenCL_C_for_context(cl_context context); 203 204 // Checks whether a particular OpenCL C version is supported by the device. 205 bool device_supports_cl_c_version(cl_device_id device, Version version); 206 207 // Poll fn every interval_ms until timeout_ms or it returns true 208 bool poll_until(unsigned timeout_ms, unsigned interval_ms, 209 std::function<bool()> fn); 210 211 // Checks whether the device supports double data types 212 bool device_supports_double(cl_device_id device); 213 214 // Checks whether the device supports half data types 215 bool device_supports_half(cl_device_id device); 216 217 #endif // _kernelHelpers_h 218