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