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 /* 44 * The below code is intended to be used at the top of kernels that appear inline in files to set line and file info for the kernel: 45 * 46 * const char *source = { 47 * INIT_OPENCL_DEBUG_INFO 48 * "__kernel void foo( int x )\n" 49 * "{\n" 50 * " ...\n" 51 * "}\n" 52 * }; 53 */ 54 #define INIT_OPENCL_DEBUG_INFO SET_OPENCL_LINE_INFO( __LINE__, __FILE__ ) 55 #define SET_OPENCL_LINE_INFO(_line, _file) "#line " STRINGIFY(_line) " " STRINGIFY(_file) "\n" 56 #ifndef STRINGIFY_VALUE 57 #define STRINGIFY_VALUE(_x) STRINGIFY(_x) 58 #endif 59 #ifndef STRINGIFY 60 #define STRINGIFY(_x) #_x 61 #endif 62 63 const int MAX_LEN_FOR_KERNEL_LIST = 20; 64 65 /* Helper that creates a single program and kernel from a single-kernel program source */ 66 extern int create_single_kernel_helper(cl_context context, 67 cl_program *outProgram, 68 cl_kernel *outKernel, 69 unsigned int numKernelLines, 70 const char **kernelProgram, 71 const char *kernelName, 72 const char *buildOptions = NULL, 73 const bool openclCXX = false); 74 75 extern int create_single_kernel_helper_with_build_options(cl_context context, 76 cl_program *outProgram, 77 cl_kernel *outKernel, 78 unsigned int numKernelLines, 79 const char **kernelProgram, 80 const char *kernelName, 81 const char *buildOptions, 82 const bool openclCXX = false); 83 84 extern int create_single_kernel_helper_create_program(cl_context context, 85 cl_program *outProgram, 86 unsigned int numKernelLines, 87 const char **kernelProgram, 88 const char *buildOptions = NULL, 89 const bool openclCXX = false); 90 91 extern int create_single_kernel_helper_create_program_for_device(cl_context context, 92 cl_device_id device, 93 cl_program *outProgram, 94 unsigned int numKernelLines, 95 const char **kernelProgram, 96 const char *buildOptions = NULL, 97 const bool openclCXX = false); 98 99 /* Creates OpenCL C++ program. This one must be used for creating OpenCL C++ program. */ 100 extern int create_openclcpp_program(cl_context context, 101 cl_program *outProgram, 102 unsigned int numKernelLines, 103 const char **kernelProgram, 104 const char *buildOptions = NULL); 105 106 /* Builds program (outProgram) and creates one kernel */ 107 int build_program_create_kernel_helper(cl_context context, 108 cl_program *outProgram, 109 cl_kernel *outKernel, 110 unsigned int numKernelLines, 111 const char **kernelProgram, 112 const char *kernelName, 113 const char *buildOptions = NULL); 114 115 /* Helper to obtain the biggest fit work group size for all the devices in a given group and for the given global thread size */ 116 extern int get_max_common_work_group_size( cl_context context, cl_kernel kernel, size_t globalThreadSize, size_t *outSize ); 117 118 /* Helper to obtain the biggest fit work group size for all the devices in a given group and for the given global thread size */ 119 extern int get_max_common_2D_work_group_size( cl_context context, cl_kernel kernel, size_t *globalThreadSize, size_t *outSizes ); 120 121 /* Helper to obtain the biggest fit work group size for all the devices in a given group and for the given global thread size */ 122 extern int get_max_common_3D_work_group_size( cl_context context, cl_kernel kernel, size_t *globalThreadSize, size_t *outSizes ); 123 124 /* Helper to obtain the biggest allowed work group size for all the devices in a given group */ 125 extern int get_max_allowed_work_group_size( cl_context context, cl_kernel kernel, size_t *outSize, size_t *outLimits ); 126 127 /* Helper to obtain the biggest allowed 1D work group size on a given device */ 128 extern int get_max_allowed_1d_work_group_size_on_device( cl_device_id device, cl_kernel kernel, size_t *outSize ); 129 130 /* Helper to determine if a device supports an image format */ 131 extern int is_image_format_supported( cl_context context, cl_mem_flags flags, cl_mem_object_type image_type, const cl_image_format *fmt ); 132 133 /* Helper to get pixel size for a pixel format */ 134 size_t get_pixel_bytes( const cl_image_format *fmt ); 135 136 /* Verify the given device supports images. */ 137 extern test_status verifyImageSupport( cl_device_id device ); 138 139 /* Checks that the given device supports images. Same as verify, but doesn't print an error */ 140 extern int checkForImageSupport( cl_device_id device ); 141 extern int checkFor3DImageSupport( cl_device_id device ); 142 143 /* Checks that a given queue property is supported on the specified device. Returns 1 if supported, 0 if not or an error. */ 144 extern int checkDeviceForQueueSupport( cl_device_id device, cl_command_queue_properties prop ); 145 146 /* Helper to obtain the min alignment for a given context, i.e the max of all min alignments for devices attached to the context*/ 147 size_t get_min_alignment(cl_context context); 148 149 /* Helper to obtain the default rounding mode for single precision computation. (Double is always CL_FP_ROUND_TO_NEAREST.) Returns 0 on error. */ 150 cl_device_fp_config get_default_rounding_mode( cl_device_id device ); 151 152 #define PASSIVE_REQUIRE_IMAGE_SUPPORT( device ) \ 153 if( checkForImageSupport( device ) ) \ 154 { \ 155 log_info( "\n\tNote: device does not support images. Skipping test...\n" ); \ 156 return 0; \ 157 } 158 159 #define PASSIVE_REQUIRE_3D_IMAGE_SUPPORT( device ) \ 160 if( checkFor3DImageSupport( device ) ) \ 161 { \ 162 log_info( "\n\tNote: device does not support 3D images. Skipping test...\n" ); \ 163 return 0; \ 164 } 165 166 #define PASSIVE_REQUIRE_FP16_SUPPORT(device) \ 167 if (!is_extension_available(device, "cl_khr_fp16")) \ 168 { \ 169 log_info("\n\tNote: device does not support fp16. Skipping test...\n"); \ 170 return 0; \ 171 } 172 173 /* Prints out the standard device header for all tests given the device to print for */ 174 extern int printDeviceHeader( cl_device_id device ); 175 176 #endif // _kernelHelpers_h 177