Merge pull request #3 from KhronosGroup/master

Update nikhiljnv/OpenCL-CTS:master with KhronosGroup/OpenCL-CTS:master

.gitignore

@@ -1,4 +0,0 @@
-# build directories
-build/
-build_lnx/
-build_win/

CMakeLists.txt

@@ -1,4 +1,4 @@
-cmake_minimum_required(VERSION 3.1)
+cmake_minimum_required(VERSION 3.5.1)
 
 set( CONFORMANCE_SUFFIX "" )
 set(CLConform_SOURCE_DIR ${CMAKE_CURRENT_SOURCE_DIR})
@@ -60,10 +60,7 @@ set(CONFORMANCE_SUFFIX "" )
 #-----------------------------------------------------------
 #Vendor Customization File can be included here to provide a way to automatically
 #build driver as a dependency of the conformance tests, or other such CMake customization
-option(USE_VENDOR_CUSTOM_FILE "Use Vendor Customization File" OFF)
+include(CMakeVendor.txt OPTIONAL)
-if(USE_VENDOR_CUSTOM_FILE)
-    include(CMakeVendor.txt OPTIONAL)
-endif(USE_VENDOR_CUSTOM_FILE)
 
 #-----------------------------------------------------------
 # Development options for OpenCL C++ tests
@@ -175,17 +172,30 @@ else()
     set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /D__SSE__")
 endif()
 
+if(MSVC)
+    # Don't warn when using standard non-secure functions.
+    add_compile_definitions(_CRT_SECURE_NO_WARNINGS)
+endif()
+
 if( WIN32 AND "${CMAKE_CXX_COMPILER_ID}" MATCHES "Intel" )
-    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Qlong-double -Qpc80 /DWIN32 /D_WINDOWS /W3 /GR /EHsc -nologo -Od -D_CRT_SECURE_NO_WARNINGS -D_CRT_NONSTDC_NO_WARNINGS -EHsc -Wall -Qdiag-disable:68,111,177,186,161,869,1028,2259,2553,181,239,265,1188 -fp:strict -fp:source")
+    set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Qlong-double -Qpc80 /DWIN32 /D_WINDOWS /W3 /GR /EHsc -nologo -Od -D_CRT_NONSTDC_NO_WARNINGS -EHsc -Wall -Qdiag-disable:68,111,177,186,161,869,1028,2259,2553,181,239,265,1188 -fp:strict -fp:source")
-    set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -Qlong-double -Qpc80 /DWIN32 /D_WINDOWS /W3 /GR /EHsc -nologo -Od -D_CRT_SECURE_NO_WARNINGS -D_CRT_NONSTDC_NO_WARNINGS -EHsc -Wall -Qdiag-disable:68,111,177,186,161,869,1028,2259,2553,181,239,265,1188 -fp:strict -fp:source")
+    set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} -Qlong-double -Qpc80 /DWIN32 /D_WINDOWS /W3 /GR /EHsc -nologo -Od -D_CRT_NONSTDC_NO_WARNINGS -EHsc -Wall -Qdiag-disable:68,111,177,186,161,869,1028,2259,2553,181,239,265,1188 -fp:strict -fp:source")
 endif()
 
 list(APPEND CLConform_LIBRARIES ${OPENCL_LIBRARIES})
 if(ANDROID)
     list(APPEND CLConform_LIBRARIES m)
-elseif(NOT WIN32)
+endif()
+if(NOT DEFINED LINK_PTHREAD)
+    if(ANDROID OR WIN32)
+        set(LINK_PTHREAD OFF)
+    else()
+        set(LINK_PTHREAD ON)
+    endif()
+endif()
+if(LINK_PTHREAD)
     list(APPEND CLConform_LIBRARIES pthread)
-endif(ANDROID)
+endif()
 
 if(APPLE)
     find_library(corefoundation CoreFoundation)

CMakeVendor.txt

@@ -1,8 +0,0 @@
-# We intentionally hardcode "_win32" to ensure backwards compatibility (to avoid breaking HAAVE)
-if(ANDROID)
-   if(ARM64_V8A)
-       set(ARCH "64")
-   else(ARM64_V8A)
-       set(ARCH "32")
-   endif(ARM64_V8A)
-endif (ANDROID)

build_android.py

@@ -1,161 +0,0 @@
-#!/usr/bin/python
-
-
-#-------------------------------------------------------------------------------#
-# android-cmake and android-ndk based build script for conformance
-#-------------------------------------------------------------------------------#
-"""
-Dependencies:
-
-1) android-ndk version android-ndk-r10d or higher is required. Further, the environment
-   variable ANDROID_NDK should be defined to point to it.
-
-2) android-cmake should be installed (else the script can install it for you). If installed,
-   the environment variable ANDROID_CMAKE should point to install location, unless it is in the current
-   working directory in which case it is picked up by default.
-
-3) CL_INCLUDE_DIR should be defined to point to CL headers. Alternately, this can be provided
-   as an input (-I)
-
-4) Path to opencl library to link against (libOpenCL.so) can be provided using -L. If this isn't
-   available the script will try to use CL_LIB_DIR_64 or CL_LIB_DIR_32 environment variables -
-   if available - to pick up the right library for the architecture being built.
-
-
-"""
-
-import os
-import sys
-import subprocess
-import argparse
-import time
-import shlex
-
-start  = time.time()
-script = os.path.basename( sys.argv[ 0 ] )
-
-def die (msg):
-    print msg
-    exit(-1)
-
-def execute (cmdline):
-    retcode = subprocess.call(cmdline)
-    if retcode != 0:
-        raise Exception("Failed to execute '%s', got %d" % (commandLine, retcode))
-
-def build(args):
-    if not (args.testDir):
-        print("building...")
-        execute("make")
-    else:
-        if os.path.exists( os.path.join(args.bld_dir, "test_conformance", args.testDir) ):
-            os.chdir( os.path.join("test_conformance",args.testDir) )
-            print("Building test: %s..." %args.testDir)
-            execute("make")
-            os.chdir(args.bld_dir)
-        else:
-            print ("Error: %s test doesn't exist" %args.testDir)
-
-
-def configure (args):
-    print("configuring...")
-    cmdline = []
-    cmdline.extend(['cmake', "-DCMAKE_TOOLCHAIN_FILE=" + os.path.join(args.android_cmake,"android.toolchain.cmake")])
-    for var in args.cmake_defs :
-        cmdline.extend([ '-D', var ])
-    cmdline.extend(['-DCL_INCLUDE_DIR=' + args.inc_dir])
-    cmdline.extend(['-DCL_LIB_DIR=' + args.lib_dir])
-    cmdline.extend(['-DANDROID_NATIVE_API_LEVEL=' + "android-21"])
-    if args.arch == "64":
-        cmdline.extend(['-DANDROID_ABI=arm64-v8a'])
-        cmdline.extend(['-DANDROID_SO_UNDEFINED=ON'])
-    cmdline.extend([args.src_dir])
-    execute(cmdline)
-
-def check_var (parser, args, name):
-    if not(args.__dict__[name]):
-        parser.error("%s needs to be defined" % name)
-
-def print_config(args):
-    print("----------CONFIGURATION--------------\n")
-    print("android_cmake: %s" % args.android_cmake)
-    print("android_ndk:   %s" % args.android_ndk)
-    print("lib_dir:       %s" % args.lib_dir)
-    print("inc_dir:       %s" % args.inc_dir)
-    if len(args.cmake_defs):
-        print("cmake options:" + "\n:".join( [ " `%s'" % dir for dir in args.cmake_defs ] ))
-    print("architecture:  %s" % args.arch)
-    print("-------------------------------------\n")
-
-def get_input():
-    yes = set(['yes','y', 'ye', ''])
-    no = set(['no','n'])
-
-    choice = raw_input().lower()
-    if choice in yes:
-        return True
-    elif choice in no:
-        return False
-    else:
-        sys.stdout.write("Please respond with 'yes' or 'no'")
-        exit()
-
-def install_android_cmake():
-    parser.print_help()
-    print "\nandroid-cmake doesn't seem to be installed - It should be provided as a) cmdline input b) environment variable $ANDROID_CMAKE or c) present in the current directory\n"
-    print "if you would like to download and install it in the current directory please enter yes\n"
-    print "if you would like to provide an environment variable($ANDROID_CMAKE) or command-line input(--android_cmake) rerun the script enter no\n"
-    print "input: "
-    if get_input():
-        print("installing android-cmake")
-        #subprocess.call(['git', 'clone', 'https://github.com/taka-no-me/android-cmake'])
-        # Use a newer fork of android-cmake which has been updated to support Clang. GCC is deprecated in newer NDKs and C11 atomics conformance doesn't build with NDK > 10.
-        subprocess.call(['git', 'clone', 'https://github.com/daewoong-jang/android-cmake'])
-        args.android_cmake = os.path.join(args.src_dir,"android-cmake")
-    else:
-        exit()
-
-try:
-    parser = argparse.ArgumentParser()
-    parser.add_argument('--android_cmake', dest='android_cmake', default=os.environ.get('ANDROID_CMAKE'), help="Path to android-cmake (can also be set using environment variable $ANDROID_CMAKE).")
-    parser.add_argument('--android_ndk', dest='android_ndk', default=os.environ.get('ANDROID_NDK'), help="Path to android-ndk (can also be set using environment variable $ANDROID_NDK).")
-    parser.add_argument('-L','--lib_dir', dest='lib_dir', default="", help="Path to libOpenCL to link against (can also be set using environment variable $CL_LIB_DIR_32 and $CL_LIB_DIR_64).")
-    parser.add_argument('-I','--include_dir', dest='inc_dir', default=os.environ.get('CL_INCLUDE_DIR'), help="Path to headers (can also be set using environment variable $CL_INCLUDE_DIR).")
-    parser.add_argument('-D', dest='cmake_defs', action='append', default=[], help="Define CMAKE variable")
-    parser.add_argument('-a','--arch', default="32", help="Architecture to build for (32 or 64)")
-    parser.add_argument('-t','--test', dest='testDir', default="", help="Builds the given test")
-
-    args = parser.parse_args()
-
-    args.src_dir = os.path.realpath(os.path.dirname( sys.argv[ 0 ]))
-
-    if not (args.android_cmake):
-        if os.path.exists(os.path.join(args.src_dir,"android-cmake")):
-            args.android_cmake = os.path.join(args.src_dir,"android-cmake")
-        else:
-            install_android_cmake()
-
-    if not (args.lib_dir):
-        lib_var_name = "CL_LIB_DIR_" + ("32" if (args.arch == "32") else "64")
-        args.lib_dir = os.environ.get(lib_var_name)
-
-    check_var(parser, args, "android_cmake")
-    check_var(parser, args, "lib_dir")
-    check_var(parser, args, "inc_dir")
-    check_var(parser, args, "android_ndk")
-
-    print_config(args)
-
-    args.bld_dir = os.path.join(args.src_dir, 'bld_android_%s' % args.arch)
-    if not os.path.exists(args.bld_dir):
-        os.makedirs(args.bld_dir)
-    os.chdir(args.bld_dir)
-
-    configure(args)
-    build(args)
-
-    sys.exit( 0 )
-
-finally:
-    finish = time.time()
-    print("Elapsed time: %.0f s." % ( finish - start ) )

build_lnx.sh

@@ -1,12 +0,0 @@
-#!/bin/sh
-
-mkdir -p build_lnx
-cd build_lnx
-cmake -G "Unix Makefiles" ../ \
-      -DKHRONOS_OFFLINE_COMPILER=<TO_SET> \
-      -DCL_LIBCLCXX_DIR=<TO_SET> \
-      -DCL_INCLUDE_DIR=<TO_SET> \
-      -DCL_LIB_DIR=<TO_SET> \
-      -DCMAKE_RUNTIME_OUTPUT_DIRECTORY=. \
-      -DOPENCL_LIBRARIES=OpenCL
-make --jobs 8

build_win.bat

@@ -1,32 +0,0 @@
-@ECHO off
-setlocal ENABLEDELAYEDEXPANSION
-
-IF DEFINED ProgramFiles(x86) SET ProgFilesDir=%ProgramFiles(x86)%
-IF NOT DEFINED ProgFilesDir SET ProgFilesDir=%ProgramFiles%
-
-rem -------------------------------- Update these to match what's on your PC ------------------------------------------------
-
-SET VCPATH="%ProgFilesDir%\Microsoft Visual Studio 14.0\Common7\IDE\devenv.com"
-
-SET PATH=%CMAKEPATH%;%PATH%
-
-rem -------------------------------------------------------------------------------------------------------------------------
-
-setlocal ENABLEDELAYEDEXPANSION
-
-call "%VS140COMNTOOLS%\vsvars32.bat"
-
-mkdir build_win
-pushd build_win
-IF NOT EXIST CLConform.sln (
-   echo "Solution file not found, running Cmake"
-   cmake -G "Visual Studio 14 2015 Win64" ..\.  -DKHRONOS_OFFLINE_COMPILER=<TO_SET> -DCL_LIBCLCXX_DIR=<TO_SET> -DCL_INCLUDE_DIR=<TO_SET> -DCL_LIB_DIR=<TO_SET> -DCMAKE_RUNTIME_OUTPUT_DIRECTORY=. -DOPENCL_LIBRARIES=OpenCL
-) else (
-   echo "Solution file found CLConform.sln "
-)
-
-echo Building CLConform.sln...
-%VCPATH% CLConform.sln /build
-
-
-GOTO:EOF

test_common/CMakeLists.txt

@@ -8,6 +8,7 @@ set(HARNESS_SOURCES
     harness/msvc9.c
     harness/crc32.cpp
     harness/errorHelpers.cpp
+    harness/featureHelpers.cpp
     harness/genericThread.cpp
     harness/imageHelpers.cpp
     harness/kernelHelpers.cpp

test_common/gl/helpers.cpp

@@ -1277,9 +1277,7 @@ void * CreateGLTexture2DMultisample( size_t width, size_t height, size_t samples
         case kUInt:
         *((unsigned int*)p) = val*0xffffffff;
           break;
-        case kHalf:
+        case kHalf: *((cl_half *)p) = convert_float_to_half(val); break;
-          *((cl_ushort*)p) = convert_float_to_half(val);
-          break;
         default:
           log_error("Test error: unexpected type enum 0x%x\n",type);
       }
@@ -1541,9 +1539,7 @@ void * CreateGLTexture2DArrayMultisample(size_t width, size_t height,
         case kUInt:
           *((unsigned int*)p) = val*0xffffffff;
           break;
-        case kHalf:
+        case kHalf: *((cl_half *)p) = convert_float_to_half(val); break;
-          *((cl_ushort*)p) = convert_float_to_half(val);
-          break;
         default:
           log_error("Test error: unexpected type enum 0x%x\n",type);
       }

test_common/harness/conversions.cpp

@@ -206,8 +206,8 @@ static Long sLowerLimits[ kNumExplicitTypes ] =
         -1,
         -128, 0, 0,
         -32768, 0, 0,
-        0xffffffff80000000LL, 0, 0,
+        (Long)0xffffffff80000000LL, 0, 0,
-        0x8000000000000000LL, 0, 0,
+        (Long)0x8000000000000000LL, 0, 0,
         0, 0 };    // Last two values aren't stored here
 
 #define BOOL_CASE(inType) \
@@ -880,7 +880,7 @@ void generate_random_data( ExplicitType type, size_t count, MTdata d, void *outD
     cl_ulong *ulongPtr;
     cl_float *floatPtr;
     cl_double *doublePtr;
-    cl_ushort *halfPtr;
+    cl_half *halfPtr;
     size_t i;
     cl_uint bits = genrand_int32(d);
     cl_uint bitsLeft = 32;

test_common/harness/deviceInfo.cpp

@@ -24,8 +24,8 @@
 
 /* Helper to return a string containing device information for the specified
  * device info parameter. */
-static std::string get_device_info_string(cl_device_id device,
+std::string get_device_info_string(cl_device_id device,
-                                          cl_device_info param_name)
+                                   cl_device_info param_name)
 {
     size_t size = 0;
     int err;

test_common/harness/deviceInfo.h

@@ -23,6 +23,11 @@
 
 #include <CL/opencl.h>
 
+/* Helper to return a string containing device information for the specified
+ * device info parameter. */
+std::string get_device_info_string(cl_device_id device,
+                                   cl_device_info param_name);
+
 /* Determines if an extension is supported by a device. */
 int is_extension_available(cl_device_id device, const char *extensionName);
 

test_common/harness/errorHelpers.cpp

@@ -354,7 +354,7 @@ static float Ulp_Error_Half_Float( float test, double reference )
     return (float) scalbn( testVal - reference, ulp_exp );
 }
 
-float Ulp_Error_Half( cl_ushort test, float reference )
+float Ulp_Error_Half(cl_half test, float reference)
 {
     return Ulp_Error_Half_Float(cl_half_to_float(test), reference);
 }

test_common/harness/errorHelpers.h

@@ -1,6 +1,6 @@
 //
 // Copyright (c) 2017 The Khronos Group Inc.
-// 
+//
 // Licensed under the Apache License, Version 2.0 (the "License");
 // you may not use this file except in compliance with the License.
 // You may obtain a copy of the License at
@@ -37,19 +37,19 @@
 #define vlog_perf(_number, _higherBetter, _numType, _format, ...) printf("Performance Number " _format " (in %s, %s): %g\n",##__VA_ARGS__, _numType,    \
                     _higherBetter?"higher is better":"lower is better" , _number)
 #ifdef _WIN32
-    #ifdef __MINGW32__
+#ifdef __MINGW32__
-        // Use __mingw_printf since it supports "%a" format specifier
+// Use __mingw_printf since it supports "%a" format specifier
-        #define vlog __mingw_printf
+#define vlog __mingw_printf
-        #define vlog_error __mingw_printf
+#define vlog_error __mingw_printf
-    #else
-        // Use home-baked function that treats "%a" as "%f"
-    static int vlog_win32(const char *format, ...);
-    #define vlog vlog_win32
-    #define vlog_error vlog_win32
-    #endif
 #else
-    #define vlog_error printf
+// Use home-baked function that treats "%a" as "%f"
-    #define vlog printf
+static int vlog_win32(const char *format, ...);
+#define vlog vlog_win32
+#define vlog_error vlog_win32
+#endif
+#else
+#define vlog_error printf
+#define vlog printf
 #endif
 
 #define ct_assert(b)          ct_assert_i(b, __LINE__)
@@ -62,16 +62,26 @@
         return TEST_FAIL;                                                      \
     }
 #define test_error(errCode,msg)    test_error_ret(errCode,msg,errCode)
-#define test_error_ret(errCode,msg,retValue)    { if( errCode != CL_SUCCESS ) { print_error( errCode, msg ); return retValue ; } }
+#define test_error_ret(errCode, msg, retValue)                                 \
+    {                                                                          \
+        auto errCodeResult = errCode;                                          \
+        if (errCodeResult != CL_SUCCESS)                                       \
+        {                                                                      \
+            print_error(errCodeResult, msg);                                   \
+            return retValue;                                                   \
+        }                                                                      \
+    }
 #define print_error(errCode,msg)    log_error( "ERROR: %s! (%s from %s:%d)\n", msg, IGetErrorString( errCode ), __FILE__, __LINE__ );
 
 #define test_missing_feature(errCode, msg) test_missing_feature_ret(errCode, msg, errCode)
-// this macro should always return CL_SUCCESS, but print the missing feature message
+// this macro should always return CL_SUCCESS, but print the missing feature
+// message
 #define test_missing_feature_ret(errCode,msg,retValue)    { if( errCode != CL_SUCCESS ) { print_missing_feature( errCode, msg ); return CL_SUCCESS ; } }
 #define print_missing_feature(errCode, msg) log_missing_feature("ERROR: Subtest %s tests a feature not supported by the device version! (from %s:%d)\n", msg, __FILE__, __LINE__ );
 
 #define test_missing_support_offline_cmpiler(errCode, msg) test_missing_support_offline_cmpiler_ret(errCode, msg, errCode)
-// this macro should always return CL_SUCCESS, but print the skip message on test not supported with offline compiler
+// this macro should always return CL_SUCCESS, but print the skip message on
+// test not supported with offline compiler
 #define test_missing_support_offline_cmpiler_ret(errCode,msg,retValue)    { if( errCode != CL_SUCCESS ) { log_info( "INFO: Subtest %s tests is not supported in offline compiler execution path! (from %s:%d)\n", msg, __FILE__, __LINE__ ); return TEST_SKIP ; } }
 
 // expected error code vs. what we got
@@ -112,9 +122,9 @@
 
 extern const char    *IGetErrorString( int clErrorCode );
 
-extern float Ulp_Error_Half( cl_ushort test, float reference );
+extern float Ulp_Error_Half(cl_half test, float reference);
-extern float Ulp_Error( float test, double reference );
+extern float Ulp_Error(float test, double reference);
-extern float Ulp_Error_Double( double test, long double reference );
+extern float Ulp_Error_Double(double test, long double reference);
 
 extern const char *GetChannelTypeName( cl_channel_type type );
 extern int IsChannelTypeSupported( cl_channel_type type );
@@ -125,7 +135,8 @@ extern const char *GetQueuePropertyName(cl_command_queue_properties properties);
 
 extern const char *GetDeviceTypeName( cl_device_type type );
 int check_functions_for_offline_compiler(const char *subtestname, cl_device_id device);
-
+cl_int OutputBuildLogs(cl_program program, cl_uint num_devices,
+                       cl_device_id *device_list);
 // NON-REENTRANT UNLESS YOU PROVIDE A BUFFER PTR (pass null to use static storage, but it's not reentrant then!)
 extern const char *GetDataVectorString( void *dataBuffer, size_t typeSize, size_t vecSize, char *buffer );
 

test_common/harness/featureHelpers.cpp

@@ -0,0 +1,75 @@
+//
+// Copyright (c) 2020 The Khronos Group Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+#include "featureHelpers.h"
+#include "errorHelpers.h"
+
+#include <assert.h>
+#include <string.h>
+
+#include <vector>
+
+int get_device_cl_c_features(cl_device_id device, OpenCLCFeatures& features)
+{
+    // Initially, all features are unsupported.
+    features = { 0 };
+
+    // The CL_DEVICE_OPENCL_C_FEATURES query does not exist pre-3.0.
+    const Version version = get_device_cl_version(device);
+    if (version < Version(3, 0))
+    {
+        return TEST_PASS;
+    }
+
+    cl_int error = CL_SUCCESS;
+
+    size_t sz = 0;
+    error = clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_FEATURES, 0, NULL, &sz);
+    test_error(error, "Unable to query CL_DEVICE_OPENCL_C_FEATURES size");
+
+    std::vector<cl_name_version> clc_features(sz / sizeof(cl_name_version));
+    error = clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_FEATURES, sz,
+                            clc_features.data(), NULL);
+    test_error(error, "Unable to query CL_DEVICE_OPENCL_C_FEATURES");
+
+#define CHECK_OPENCL_C_FEATURE(_feature)                                       \
+    if (strcmp(clc_feature.name, #_feature) == 0)                              \
+    {                                                                          \
+        features.supports##_feature = true;                                    \
+    }
+
+    for (const auto& clc_feature : clc_features)
+    {
+        CHECK_OPENCL_C_FEATURE(__opencl_c_3d_image_writes);
+        CHECK_OPENCL_C_FEATURE(__opencl_c_atomic_order_acq_rel);
+        CHECK_OPENCL_C_FEATURE(__opencl_c_atomic_order_seq_cst);
+        CHECK_OPENCL_C_FEATURE(__opencl_c_atomic_scope_device);
+        CHECK_OPENCL_C_FEATURE(__opencl_c_atomic_scope_all_devices);
+        CHECK_OPENCL_C_FEATURE(__opencl_c_device_enqueue);
+        CHECK_OPENCL_C_FEATURE(__opencl_c_generic_address_space);
+        CHECK_OPENCL_C_FEATURE(__opencl_c_fp64);
+        CHECK_OPENCL_C_FEATURE(__opencl_c_images);
+        CHECK_OPENCL_C_FEATURE(__opencl_c_int64);
+        CHECK_OPENCL_C_FEATURE(__opencl_c_pipes);
+        CHECK_OPENCL_C_FEATURE(__opencl_c_program_scope_global_variables);
+        CHECK_OPENCL_C_FEATURE(__opencl_c_read_write_images);
+        CHECK_OPENCL_C_FEATURE(__opencl_c_subgroups);
+        CHECK_OPENCL_C_FEATURE(__opencl_c_work_group_collective_functions);
+    }
+
+#undef CHECK_OPENCL_C_FEATURE
+
+    return TEST_PASS;
+}

test_common/harness/featureHelpers.h

@@ -0,0 +1,43 @@
+//
+// Copyright (c) 2020 The Khronos Group Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+#ifndef _featureHelpers_h
+#define _featureHelpers_h
+
+#include "compat.h"
+#include "testHarness.h"
+
+struct OpenCLCFeatures
+{
+    bool supports__opencl_c_3d_image_writes;
+    bool supports__opencl_c_atomic_order_acq_rel;
+    bool supports__opencl_c_atomic_order_seq_cst;
+    bool supports__opencl_c_atomic_scope_device;
+    bool supports__opencl_c_atomic_scope_all_devices;
+    bool supports__opencl_c_device_enqueue;
+    bool supports__opencl_c_generic_address_space;
+    bool supports__opencl_c_fp64;
+    bool supports__opencl_c_images;
+    bool supports__opencl_c_int64;
+    bool supports__opencl_c_pipes;
+    bool supports__opencl_c_program_scope_global_variables;
+    bool supports__opencl_c_read_write_images;
+    bool supports__opencl_c_subgroups;
+    bool supports__opencl_c_work_group_collective_functions;
+};
+
+int get_device_cl_c_features(cl_device_id device, OpenCLCFeatures& features);
+
+#endif // _featureHelpers_h

test_common/harness/imageHelpers.cpp

@@ -925,7 +925,7 @@ int get_format_min_int( cl_image_format *format )
     }
 }
 
-cl_ushort convert_float_to_half( float f )
+cl_half convert_float_to_half(float f)
 {
     switch( gFloatToHalfRoundingMode )
     {
@@ -1281,10 +1281,9 @@ void read_image_pixel_float( void *imageData, image_descriptor *imageInfo,
             break;
         }
 
-        case CL_HALF_FLOAT:
+        case CL_HALF_FLOAT: {
-        {
+            cl_half *dPtr = (cl_half *)ptr;
-            cl_ushort *dPtr = (cl_ushort *)ptr;
+            for (i = 0; i < channelCount; i++)
-            for( i = 0; i < channelCount; i++ )
                 tempData[i] = cl_half_to_float(dPtr[i]);
             break;
         }
@@ -2397,9 +2396,8 @@ void pack_image_pixel( float *srcVector, const cl_image_format *imageFormat, voi
     size_t channelCount = get_format_channel_count( imageFormat );
     switch( imageFormat->image_channel_data_type )
     {
-        case CL_HALF_FLOAT:
+        case CL_HALF_FLOAT: {
-        {
+            cl_half *ptr = (cl_half *)outData;
-            cl_ushort *ptr = (cl_ushort *)outData;
 
             switch( gFloatToHalfRoundingMode )
             {
@@ -2569,9 +2567,8 @@ void pack_image_pixel_error( const float *srcVector, const cl_image_format *imag
     size_t channelCount = get_format_channel_count( imageFormat );
     switch( imageFormat->image_channel_data_type )
     {
-        case CL_HALF_FLOAT:
+        case CL_HALF_FLOAT: {
-        {
+            const cl_half *ptr = (const cl_half *)results;
-            const cl_ushort *ptr = (const cl_ushort *)results;
 
             for( unsigned int i = 0; i < channelCount; i++ )
                 errors[i] = Ulp_Error_Half( ptr[i], srcVector[i] );
@@ -2838,25 +2835,28 @@ int  DetectFloatToHalfRoundingMode( cl_command_queue q )  // Returns CL_SUCCESS
             return err;
         }
 
-    // read the results
+        // read the results
-        cl_ushort outBuf[count*4];
+        cl_half outBuf[count * 4];
-        memset( outBuf, -1, sizeof( outBuf ) );
+        memset(outBuf, -1, sizeof(outBuf));
-        size_t origin[3] = {0,0,0};
+        size_t origin[3] = { 0, 0, 0 };
-        size_t region[3] = {count,1,1};
+        size_t region[3] = { count, 1, 1 };
-        err = clEnqueueReadImage( q, outImage, CL_TRUE, origin, region, 0, 0, outBuf, 0, NULL, NULL );
+        err = clEnqueueReadImage(q, outImage, CL_TRUE, origin, region, 0, 0,
-        if( err )
+                                 outBuf, 0, NULL, NULL);
+        if (err)
         {
-            log_error( "Error: could not read output image in DetectFloatToHalfRoundingMode (%d)", err );
+            log_error("Error: could not read output image in "
-            clReleaseMemObject( inBuf );
+                      "DetectFloatToHalfRoundingMode (%d)",
-            clReleaseMemObject( outImage );
+                      err);
-            clReleaseKernel( k );
+            clReleaseMemObject(inBuf);
+            clReleaseMemObject(outImage);
+            clReleaseKernel(k);
             return err;
         }
 
-    // Generate our list of reference results
+        // Generate our list of reference results
-        cl_ushort rte_ref[count*4];
+        cl_half rte_ref[count * 4];
-        cl_ushort rtz_ref[count*4];
+        cl_half rtz_ref[count * 4];
-        for( size_t i = 0; i < 4 * count; i++ )
+        for (size_t i = 0; i < 4 * count; i++)
         {
             rte_ref[i] = cl_half_from_float(inp[i], CL_HALF_RTE);
             rtz_ref[i] = cl_half_from_float(inp[i], CL_HALF_RTZ);
@@ -3462,150 +3462,199 @@ bool find_format( cl_image_format *formatList, unsigned int numFormats, cl_image
     return false;
 }
 
-void build_required_image_formats(cl_mem_flags flags,
+void build_required_image_formats(
-                                  cl_mem_object_type image_type,
+    cl_mem_flags flags, cl_mem_object_type image_type, cl_device_id device,
-                                  cl_device_id device,
+    std::vector<cl_image_format> &formatsToSupport)
-                                  std::vector<cl_image_format>& formatsToSupport)
 {
-	Version version = get_device_cl_version(device);
+    formatsToSupport.clear();
 
-	formatsToSupport.clear();
+    // Minimum list of supported image formats for reading or writing (embedded
+    // profile)
+    static std::vector<cl_image_format> embeddedProfile_readOrWrite{
+        // clang-format off
+        { CL_RGBA, CL_UNORM_INT8 },
+        { CL_RGBA, CL_UNORM_INT16 },
+        { CL_RGBA, CL_SIGNED_INT8 },
+        { CL_RGBA, CL_SIGNED_INT16 },
+        { CL_RGBA, CL_SIGNED_INT32 },
+        { CL_RGBA, CL_UNSIGNED_INT8 },
+        { CL_RGBA, CL_UNSIGNED_INT16 },
+        { CL_RGBA, CL_UNSIGNED_INT32 },
+        { CL_RGBA, CL_HALF_FLOAT },
+        { CL_RGBA, CL_FLOAT },
+        // clang-format on
+    };
 
-	// Required embedded formats.
+    // Minimum list of required image formats for reading or writing
-	static std::vector<cl_image_format> embeddedProfReadOrWriteFormats
+    // num_channels, for all image types.
-	{
+    static std::vector<cl_image_format> fullProfile_readOrWrite{
-		{ CL_RGBA, CL_UNORM_INT8 },
+        // clang-format off
-		{ CL_RGBA, CL_UNORM_INT16 },
+        { CL_RGBA, CL_UNORM_INT8 },
-		{ CL_RGBA, CL_SIGNED_INT8 },
+        { CL_RGBA, CL_UNORM_INT16 },
-		{ CL_RGBA, CL_SIGNED_INT16 },
+        { CL_RGBA, CL_SIGNED_INT8 },
-		{ CL_RGBA, CL_SIGNED_INT32 },
+        { CL_RGBA, CL_SIGNED_INT16 },
-		{ CL_RGBA, CL_UNSIGNED_INT8 },
+        { CL_RGBA, CL_SIGNED_INT32 },
-		{ CL_RGBA, CL_UNSIGNED_INT16 },
+        { CL_RGBA, CL_UNSIGNED_INT8 },
-		{ CL_RGBA, CL_UNSIGNED_INT32 },
+        { CL_RGBA, CL_UNSIGNED_INT16 },
-		{ CL_RGBA, CL_HALF_FLOAT },
+        { CL_RGBA, CL_UNSIGNED_INT32 },
-		{ CL_RGBA, CL_FLOAT },
+        { CL_RGBA, CL_HALF_FLOAT },
-	};
+        { CL_RGBA, CL_FLOAT },
+        { CL_BGRA, CL_UNORM_INT8 },
+        // clang-format on
+    };
 
-	/*
+    // Minimum list of supported image formats for reading or writing
-		Required full profile formats.
+    // (OpenCL 2.0, 2.1, or 2.2), for all image types.
-		This array does not contain any full profile
+    static std::vector<cl_image_format> fullProfile_2x_readOrWrite{
-		formats that have restrictions on when they
+        // clang-format off
-		are required.
+        { CL_R, CL_UNORM_INT8 },
-	*/
+        { CL_R, CL_UNORM_INT16 },
-	static std::vector<cl_image_format> fullProfReadOrWriteFormats
+        { CL_R, CL_SNORM_INT8 },
-	{
+        { CL_R, CL_SNORM_INT16 },
-		{ CL_RGBA, CL_UNORM_INT8 },
+        { CL_R, CL_SIGNED_INT8 },
-		{ CL_RGBA, CL_UNORM_INT16 },
+        { CL_R, CL_SIGNED_INT16 },
-		{ CL_RGBA, CL_SIGNED_INT8 },
+        { CL_R, CL_SIGNED_INT32 },
-		{ CL_RGBA, CL_SIGNED_INT16 },
+        { CL_R, CL_UNSIGNED_INT8 },
-		{ CL_RGBA, CL_SIGNED_INT32 },
+        { CL_R, CL_UNSIGNED_INT16 },
-		{ CL_RGBA, CL_UNSIGNED_INT8 },
+        { CL_R, CL_UNSIGNED_INT32 },
-		{ CL_RGBA, CL_UNSIGNED_INT16 },
+        { CL_R, CL_HALF_FLOAT },
-		{ CL_RGBA, CL_UNSIGNED_INT32 },
+        { CL_R, CL_FLOAT },
-		{ CL_RGBA, CL_HALF_FLOAT },
+        { CL_RG, CL_UNORM_INT8 },
-		{ CL_RGBA, CL_FLOAT },
+        { CL_RG, CL_UNORM_INT16 },
-		{ CL_BGRA, CL_UNORM_INT8 },
+        { CL_RG, CL_SNORM_INT8 },
-	};
+        { CL_RG, CL_SNORM_INT16 },
+        { CL_RG, CL_SIGNED_INT8 },
+        { CL_RG, CL_SIGNED_INT16 },
+        { CL_RG, CL_SIGNED_INT32 },
+        { CL_RG, CL_UNSIGNED_INT8 },
+        { CL_RG, CL_UNSIGNED_INT16 },
+        { CL_RG, CL_UNSIGNED_INT32 },
+        { CL_RG, CL_HALF_FLOAT },
+        { CL_RG, CL_FLOAT },
+        { CL_RGBA, CL_UNORM_INT8 },
+        { CL_RGBA, CL_UNORM_INT16 },
+        { CL_RGBA, CL_SNORM_INT8 },
+        { CL_RGBA, CL_SNORM_INT16 },
+        { CL_RGBA, CL_SIGNED_INT8 },
+        { CL_RGBA, CL_SIGNED_INT16 },
+        { CL_RGBA, CL_SIGNED_INT32 },
+        { CL_RGBA, CL_UNSIGNED_INT8 },
+        { CL_RGBA, CL_UNSIGNED_INT16 },
+        { CL_RGBA, CL_UNSIGNED_INT32 },
+        { CL_RGBA, CL_HALF_FLOAT },
+        { CL_RGBA, CL_FLOAT },
+        { CL_BGRA, CL_UNORM_INT8 },
+        // clang-format on
+    };
 
-	/*
+    // Conditional addition to the 2x readOrWrite table:
-		Required full profile formats specifically for 2.x.
+    // Support for the CL_DEPTH image channel order is required only for 2D
-		This array does not contain any full profile
+    // images and 2D image arrays.
-		formats that have restrictions on when they
+    static std::vector<cl_image_format> fullProfile_2x_readOrWrite_Depth{
-		are required.
+        // clang-format off
-	*/
+        { CL_DEPTH, CL_UNORM_INT16 },
-	static std::vector<cl_image_format> fullProf2XReadOrWriteFormats
+        { CL_DEPTH, CL_FLOAT },
-	{
+        // clang-format on
-		{ CL_R, CL_UNORM_INT8 },
+    };
-		{ CL_R, CL_UNORM_INT16 },
-		{ CL_R, CL_SNORM_INT8 },
-		{ CL_R, CL_SNORM_INT16 },
-		{ CL_R, CL_SIGNED_INT8 },
-		{ CL_R, CL_SIGNED_INT16 },
-		{ CL_R, CL_SIGNED_INT32 },
-		{ CL_R, CL_UNSIGNED_INT8 },
-		{ CL_R, CL_UNSIGNED_INT16 },
-		{ CL_R, CL_UNSIGNED_INT32 },
-		{ CL_R, CL_HALF_FLOAT },
-		{ CL_R, CL_FLOAT },
-		{ CL_RG, CL_UNORM_INT8 },
-		{ CL_RG, CL_UNORM_INT16 },
-		{ CL_RG, CL_SNORM_INT8 },
-		{ CL_RG, CL_SNORM_INT16 },
-		{ CL_RG, CL_SIGNED_INT8 },
-		{ CL_RG, CL_SIGNED_INT16 },
-		{ CL_RG, CL_SIGNED_INT32 },
-		{ CL_RG, CL_UNSIGNED_INT8 },
-		{ CL_RG, CL_UNSIGNED_INT16 },
-		{ CL_RG, CL_UNSIGNED_INT32 },
-		{ CL_RG, CL_HALF_FLOAT },
-		{ CL_RG, CL_FLOAT },
-		{ CL_RGBA, CL_SNORM_INT8 },
-		{ CL_RGBA, CL_SNORM_INT16 },
-	};
 
-	/*
+    // Conditional addition to the 2x readOrWrite table:
-		Required full profile formats for CL_DEPTH
+    // Support for reading from the CL_sRGBA image channel order is optional for
-		(specifically 2.x).
+    // 1D image buffers. Support for writing to the CL_sRGBA image channel order
-		There are cases whereby the format isn't required.
+    // is optional for all image types.
-	*/
+    static std::vector<cl_image_format> fullProfile_2x_readOrWrite_srgb{
-	static std::vector<cl_image_format> fullProf2XReadOrWriteDepthFormats
+        { CL_sRGBA, CL_UNORM_INT8 },
-	{
+    };
-		{ CL_DEPTH, CL_UNORM_INT16 },
-		{ CL_DEPTH, CL_FLOAT },
-	};
 
-	/*
+    // Minimum list of required image formats for reading and writing.
-		Required full profile formats for CL_sRGB
+    static std::vector<cl_image_format> fullProfile_readAndWrite{
-		(specifically 2.x).
+        // clang-format off
-		There are cases whereby the format isn't required.
+        { CL_R, CL_UNORM_INT8 },
-	*/
+        { CL_R, CL_SIGNED_INT8 },
-	static std::vector<cl_image_format> fullProf2XSRGBFormats
+        { CL_R, CL_SIGNED_INT16 },
-	{
+        { CL_R, CL_SIGNED_INT32 },
-		{ CL_sRGBA, CL_UNORM_INT8 },
+        { CL_R, CL_UNSIGNED_INT8 },
-	};
+        { CL_R, CL_UNSIGNED_INT16 },
+        { CL_R, CL_UNSIGNED_INT32 },
+        { CL_R, CL_HALF_FLOAT },
+        { CL_R, CL_FLOAT },
+        { CL_RGBA, CL_UNORM_INT8 },
+        { CL_RGBA, CL_SIGNED_INT8 },
+        { CL_RGBA, CL_SIGNED_INT16 },
+        { CL_RGBA, CL_SIGNED_INT32 },
+        { CL_RGBA, CL_UNSIGNED_INT8 },
+        { CL_RGBA, CL_UNSIGNED_INT16 },
+        { CL_RGBA, CL_UNSIGNED_INT32 },
+        { CL_RGBA, CL_HALF_FLOAT },
+        { CL_RGBA, CL_FLOAT },
+        // clang-format on
+    };
 
-	// Embedded profile
+    // Embedded profile
-	if (gIsEmbedded)
+    if (gIsEmbedded)
-	{
+    {
-		copy(embeddedProfReadOrWriteFormats.begin(),
+        copy(embeddedProfile_readOrWrite.begin(),
-		     embeddedProfReadOrWriteFormats.end(),
+             embeddedProfile_readOrWrite.end(),
-		     back_inserter(formatsToSupport));
+             back_inserter(formatsToSupport));
-	}
+    }
-	// Full profile
+    // Full profile
-	else
+    else
-	{
+    {
-		copy(fullProfReadOrWriteFormats.begin(),
+        Version version = get_device_cl_version(device);
-		     fullProfReadOrWriteFormats.end(),
+        if (version < Version(2, 0) || version >= Version(3, 0))
-		     back_inserter(formatsToSupport));
+        {
-	}
+            // Full profile, OpenCL 1.2 or 3.0.
+            if (flags & CL_MEM_KERNEL_READ_AND_WRITE)
+            {
+                // Note: assumes that read-write images are supported!
+                copy(fullProfile_readAndWrite.begin(),
+                     fullProfile_readAndWrite.end(),
+                     back_inserter(formatsToSupport));
+            }
+            else
+            {
+                copy(fullProfile_readOrWrite.begin(),
+                     fullProfile_readOrWrite.end(),
+                     back_inserter(formatsToSupport));
+            }
+        }
+        else
+        {
+            // Full profile, OpenCL 2.0, 2.1, 2.2.
+            if (flags & CL_MEM_KERNEL_READ_AND_WRITE)
+            {
+                copy(fullProfile_readAndWrite.begin(),
+                     fullProfile_readAndWrite.end(),
+                     back_inserter(formatsToSupport));
+            }
+            else
+            {
+                copy(fullProfile_2x_readOrWrite.begin(),
+                     fullProfile_2x_readOrWrite.end(),
+                     back_inserter(formatsToSupport));
 
-	// Full profile, OpenCL 2.0, 2.1, 2.2
+                // Support for the CL_DEPTH image channel order is required only
-	if (!gIsEmbedded && version >= Version(2, 0) && version <= Version(2, 2))
+                // for 2D images and 2D image arrays.
-	{
+                if (image_type == CL_MEM_OBJECT_IMAGE2D
-		copy(fullProf2XReadOrWriteFormats.begin(),
+                    || image_type == CL_MEM_OBJECT_IMAGE2D_ARRAY)
-		     fullProf2XReadOrWriteFormats.end(),
+                {
-		     back_inserter(formatsToSupport));
+                    copy(fullProfile_2x_readOrWrite_Depth.begin(),
+                         fullProfile_2x_readOrWrite_Depth.end(),
+                         back_inserter(formatsToSupport));
+                }
 
-		// Depth images are only required for 2DArray and 2D images
+                // Support for reading from the CL_sRGBA image channel order is
-		if (image_type == CL_MEM_OBJECT_IMAGE2D_ARRAY || image_type == CL_MEM_OBJECT_IMAGE2D)
+                // optional for 1D image buffers. Support for writing to the
-		{
+                // CL_sRGBA image channel order is optional for all image types.
-			copy(fullProf2XReadOrWriteDepthFormats.begin(),
+                if (image_type != CL_MEM_OBJECT_IMAGE1D_BUFFER
-			     fullProf2XReadOrWriteDepthFormats.end(),
+                    && flags == CL_MEM_READ_ONLY)
-			     back_inserter(formatsToSupport));
+                {
-		}
+                    copy(fullProfile_2x_readOrWrite_srgb.begin(),
-
+                         fullProfile_2x_readOrWrite_srgb.end(),
-		// sRGB is not required for 1DImage Buffers
+                         back_inserter(formatsToSupport));
-		if (image_type != CL_MEM_OBJECT_IMAGE1D_BUFFER)
+                }
-		{
+            }
-			// sRGB is only required for reading
+        }
-			if (flags == CL_MEM_READ_ONLY)
+    }
-			{
-				copy(fullProf2XSRGBFormats.begin(),
-				     fullProf2XSRGBFormats.end(),
-				     back_inserter(formatsToSupport));
-			}
-		}
-	}
 }
 
 bool is_image_format_required(cl_image_format format,
@@ -3734,3 +3783,17 @@ size_t compute_mip_level_offset( image_descriptor * imageInfo , size_t lod)
   }
   return retOffset;
 }
+
+const char *convert_image_type_to_string(cl_mem_object_type image_type)
+{
+    switch (image_type)
+    {
+        case CL_MEM_OBJECT_IMAGE1D: return "1D";
+        case CL_MEM_OBJECT_IMAGE2D: return "2D";
+        case CL_MEM_OBJECT_IMAGE3D: return "3D";
+        case CL_MEM_OBJECT_IMAGE1D_ARRAY: return "1D array";
+        case CL_MEM_OBJECT_IMAGE2D_ARRAY: return "2D array";
+        case CL_MEM_OBJECT_IMAGE1D_BUFFER: return "1D image buffer";
+        default: return "unrecognized object type";
+    }
+}

test_common/harness/imageHelpers.h

@@ -275,10 +275,9 @@ template <class T> void read_image_pixel( void *imageData, image_descriptor *ima
             break;
         }
 
-        case CL_HALF_FLOAT:
+        case CL_HALF_FLOAT: {
-        {
+            cl_half *dPtr = (cl_half *)ptr;
-            cl_ushort *dPtr = (cl_ushort *)ptr;
+            for (i = 0; i < get_format_channel_count(format); i++)
-            for( i = 0; i < get_format_channel_count( format ); i++ )
                 tempData[i] = (T)cl_half_to_float(dPtr[i]);
             break;
         }
@@ -639,18 +638,22 @@ protected:
     size_t    mVecSize;
 };
 
-extern cl_ushort convert_float_to_half(float f);
+extern cl_half convert_float_to_half(float f);
-extern int  DetectFloatToHalfRoundingMode( cl_command_queue );  // Returns CL_SUCCESS on success
+extern int DetectFloatToHalfRoundingMode(
+    cl_command_queue); // Returns CL_SUCCESS on success
 
 // sign bit: don't care, exponent: maximum value, significand: non-zero
-static int inline is_half_nan( cl_ushort half ){ return ( half & 0x7fff ) > 0x7c00; }
+static int inline is_half_nan(cl_half half) { return (half & 0x7fff) > 0x7c00; }
 
 // sign bit: don't care, exponent: zero, significand: non-zero
-static int inline is_half_denorm( cl_ushort half ){ return IsHalfSubnormal( half ); }
+static int inline is_half_denorm(cl_half half) { return IsHalfSubnormal(half); }
 
 // sign bit: don't care, exponent: zero, significand: zero
-static int inline is_half_zero( cl_ushort half ){ return ( half & 0x7fff ) == 0; }
+static int inline is_half_zero(cl_half half) { return (half & 0x7fff) == 0; }
 
 extern double sRGBmap(float fc);
 
+extern const char *convert_image_type_to_string(cl_mem_object_type imageType);
+
+
 #endif // _imageHelpers_h

test_common/harness/kernelHelpers.cpp

@@ -1616,48 +1616,22 @@ int printDeviceHeader( cl_device_id device )
              deviceName, deviceVendor, deviceVersion, ( error == CL_SUCCESS ) ? ", CL C Version = " : "",
              ( error == CL_SUCCESS ) ? cLangVersion : "" );
 
+    auto version = get_device_cl_version(device);
+    if (version >= Version(3, 0))
+    {
+        auto ctsVersion = get_device_info_string(
+            device, CL_DEVICE_LATEST_CONFORMANCE_VERSION_PASSED);
+        log_info("Device latest conformance version passed: %s\n",
+                 ctsVersion.c_str());
+    }
+
     return CL_SUCCESS;
 }
 
 Version get_device_cl_c_version(cl_device_id device)
 {
-    // Get the device OpenCL version.
     auto device_cl_version = get_device_cl_version(device);
 
-    // If the device version >= 3.0 it must support the
-    // CL_DEVICE_OPENCL_C_ALL_VERSIONS query from which we can extract the most
-    // recent CL C version supported by the device.
-    if (device_cl_version >= Version{ 3, 0 })
-    {
-        size_t opencl_c_all_versions_size_in_bytes{};
-        auto error =
-            clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, 0, nullptr,
-                            &opencl_c_all_versions_size_in_bytes);
-        test_error_ret(
-            error, "clGetDeviceInfo failed for CL_DEVICE_OPENCL_C_ALL_VERSIONS",
-            (Version{ -1, 0 }));
-        std::vector<cl_name_version> name_versions(
-            opencl_c_all_versions_size_in_bytes / sizeof(cl_name_version));
-        error = clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_ALL_VERSIONS,
-                                opencl_c_all_versions_size_in_bytes,
-                                name_versions.data(), nullptr);
-        test_error_ret(
-            error, "clGetDeviceInfo failed for CL_DEVICE_OPENCL_C_ALL_VERSIONS",
-            (Version{ -1, 0 }));
-
-        Version max_supported_cl_c_version{};
-        for (const auto &name_version : name_versions)
-        {
-            Version current_version{ CL_VERSION_MAJOR(name_version.version),
-                                     CL_VERSION_MINOR(name_version.version) };
-            max_supported_cl_c_version =
-                (current_version > max_supported_cl_c_version)
-                ? current_version
-                : max_supported_cl_c_version;
-        }
-        return max_supported_cl_c_version;
-    }
-
     // The second special case is OpenCL-1.0 where CL_DEVICE_OPENCL_C_VERSION
     // did not exist, but since this is just the first version we can
     // return 1.0.
@@ -1693,6 +1667,47 @@ Version get_device_cl_c_version(cl_device_id device)
     return Version{ major - '0', minor - '0' };
 }
 
+Version get_device_latest_cl_c_version(cl_device_id device)
+{
+    auto device_cl_version = get_device_cl_version(device);
+
+    // If the device version >= 3.0 it must support the
+    // CL_DEVICE_OPENCL_C_ALL_VERSIONS query from which we can extract the most
+    // recent CL C version supported by the device.
+    if (device_cl_version >= Version{ 3, 0 })
+    {
+        size_t opencl_c_all_versions_size_in_bytes{};
+        auto error =
+            clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, 0, nullptr,
+                            &opencl_c_all_versions_size_in_bytes);
+        test_error_ret(
+            error, "clGetDeviceInfo failed for CL_DEVICE_OPENCL_C_ALL_VERSIONS",
+            (Version{ -1, 0 }));
+        std::vector<cl_name_version> name_versions(
+            opencl_c_all_versions_size_in_bytes / sizeof(cl_name_version));
+        error = clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_ALL_VERSIONS,
+                                opencl_c_all_versions_size_in_bytes,
+                                name_versions.data(), nullptr);
+        test_error_ret(
+            error, "clGetDeviceInfo failed for CL_DEVICE_OPENCL_C_ALL_VERSIONS",
+            (Version{ -1, 0 }));
+
+        Version max_supported_cl_c_version{};
+        for (const auto &name_version : name_versions)
+        {
+            Version current_version{ CL_VERSION_MAJOR(name_version.version),
+                                     CL_VERSION_MINOR(name_version.version) };
+            max_supported_cl_c_version =
+                (current_version > max_supported_cl_c_version)
+                ? current_version
+                : max_supported_cl_c_version;
+        }
+        return max_supported_cl_c_version;
+    }
+
+    return get_device_cl_c_version(device);
+}
+
 Version get_max_OpenCL_C_for_context(cl_context context)
 {
     // Get all the devices in the context and find the maximum
@@ -1706,10 +1721,11 @@ Version get_max_OpenCL_C_for_context(cl_context context)
                                       / sizeof(cl_device_id));
     error = clGetContextInfo(context, CL_CONTEXT_DEVICES, devices_size_in_bytes,
                              devices.data(), nullptr);
-    auto current_version = get_device_cl_c_version(devices[0]);
+    auto current_version = get_device_latest_cl_c_version(devices[0]);
     std::for_each(std::next(devices.begin()), devices.end(),
                   [&current_version](cl_device_id device) {
-                      auto device_version = get_device_cl_c_version(device);
+                      auto device_version =
+                          get_device_latest_cl_c_version(device);
                       // OpenCL 3.0 is not backwards compatible with 2.0.
                       // If we have 3.0 and 2.0 in the same driver we
                       // use 1.2.
@@ -1731,6 +1747,50 @@ Version get_max_OpenCL_C_for_context(cl_context context)
     return current_version;
 }
 
+bool device_supports_cl_c_version(cl_device_id device, Version version)
+{
+    auto device_cl_version = get_device_cl_version(device);
+
+    // In general, a device does not support an OpenCL C version if it is <=
+    // CL_DEVICE_OPENCL_C_VERSION AND it does not appear in the
+    // CL_DEVICE_OPENCL_C_ALL_VERSIONS query.
+
+    // If the device version >= 3.0 it must support the
+    // CL_DEVICE_OPENCL_C_ALL_VERSIONS query, and the version of OpenCL C being
+    // used must appear in the query result if it's <=
+    // CL_DEVICE_OPENCL_C_VERSION.
+    if (device_cl_version >= Version{ 3, 0 })
+    {
+        size_t opencl_c_all_versions_size_in_bytes{};
+        auto error =
+            clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_ALL_VERSIONS, 0, nullptr,
+                            &opencl_c_all_versions_size_in_bytes);
+        test_error_ret(
+            error, "clGetDeviceInfo failed for CL_DEVICE_OPENCL_C_ALL_VERSIONS",
+            (false));
+        std::vector<cl_name_version> name_versions(
+            opencl_c_all_versions_size_in_bytes / sizeof(cl_name_version));
+        error = clGetDeviceInfo(device, CL_DEVICE_OPENCL_C_ALL_VERSIONS,
+                                opencl_c_all_versions_size_in_bytes,
+                                name_versions.data(), nullptr);
+        test_error_ret(
+            error, "clGetDeviceInfo failed for CL_DEVICE_OPENCL_C_ALL_VERSIONS",
+            (false));
+
+        for (const auto &name_version : name_versions)
+        {
+            Version current_version{ CL_VERSION_MAJOR(name_version.version),
+                                     CL_VERSION_MINOR(name_version.version) };
+            if (current_version == version)
+            {
+                return true;
+            }
+        }
+    }
+
+    return version <= get_device_cl_c_version(device);
+}
+
 bool poll_until(unsigned timeout_ms, unsigned interval_ms,
                 std::function<bool()> fn)
 {

test_common/harness/kernelHelpers.h

@@ -152,38 +152,48 @@ size_t get_min_alignment(cl_context context);
 /* Helper to obtain the default rounding mode for single precision computation. (Double is always CL_FP_ROUND_TO_NEAREST.) Returns 0 on error. */
 cl_device_fp_config get_default_rounding_mode( cl_device_id device );
 
-#define PASSIVE_REQUIRE_IMAGE_SUPPORT( device )    \
+#define PASSIVE_REQUIRE_IMAGE_SUPPORT(device)                                  \
-    if( checkForImageSupport( device ) )    \
+    if (checkForImageSupport(device))                                          \
-    {    \
+    {                                                                          \
-        log_info( "\n\tNote: device does not support images. Skipping test...\n" );    \
+        log_info(                                                              \
-        return 0;    \
+            "\n\tNote: device does not support images. Skipping test...\n");   \
+        return TEST_SKIPPED_ITSELF;                                            \
     }
 
-#define PASSIVE_REQUIRE_3D_IMAGE_SUPPORT( device )    \
+#define PASSIVE_REQUIRE_3D_IMAGE_SUPPORT(device)                               \
-    if( checkFor3DImageSupport( device ) )    \
+    if (checkFor3DImageSupport(device))                                        \
-    {    \
+    {                                                                          \
-        log_info( "\n\tNote: device does not support 3D images. Skipping test...\n" );    \
+        log_info("\n\tNote: device does not support 3D images. Skipping "      \
-        return 0;    \
+                 "test...\n");                                                 \
+        return TEST_SKIPPED_ITSELF;                                            \
     }
 
-#define PASSIVE_REQUIRE_FP16_SUPPORT(device)                            \
+#define PASSIVE_REQUIRE_FP16_SUPPORT(device)                                   \
-    if (!is_extension_available(device, "cl_khr_fp16"))                 \
+    if (!is_extension_available(device, "cl_khr_fp16"))                        \
-    {                                                                   \
+    {                                                                          \
-        log_info("\n\tNote: device does not support fp16. Skipping test...\n"); \
+        log_info(                                                              \
-        return 0;                                                       \
+            "\n\tNote: device does not support fp16. Skipping test...\n");     \
+        return TEST_SKIPPED_ITSELF;                                            \
     }
 
 /* Prints out the standard device header for all tests given the device to print for */
 extern int printDeviceHeader( cl_device_id device );
 
+// Execute the CL_DEVICE_OPENCL_C_VERSION query and return the OpenCL C version
+// is supported by the device.
+Version get_device_cl_c_version(cl_device_id device);
+
 // Gets the latest (potentially non-backward compatible) OpenCL C version
 // supported by the device.
-Version get_device_cl_c_version(cl_device_id device);
+Version get_device_latest_cl_c_version(cl_device_id device);
 
 // Gets the maximum universally supported OpenCL C version in a context, i.e.
 // the OpenCL C version supported by all devices in a context.
 Version get_max_OpenCL_C_for_context(cl_context context);
 
+// Checks whether a particular OpenCL C version is supported by the device.
+bool device_supports_cl_c_version(cl_device_id device, Version version);
+
 // Poll fn every interval_ms until timeout_ms or it returns true
 bool poll_until(unsigned timeout_ms, unsigned interval_ms,
                 std::function<bool()> fn);

test_common/harness/testHarness.cpp

@@ -15,6 +15,7 @@
 //
 #include "testHarness.h"
 #include "compat.h"
+#include <algorithm>
 #include <stdio.h>
 #include <stdlib.h>
 #include <string.h>
@@ -54,8 +55,6 @@ cl_uint gReSeed = 0;
 int     gFlushDenormsToZero = 0;
 int     gInfNanSupport = 1;
 int     gIsEmbedded = 0;
-int     gIsOpenCL_C_1_0_Device = 0;
-int     gIsOpenCL_1_0_Device = 0;
 int     gHasLong = 1;
 bool gCoreILProgram = true;
 
@@ -403,27 +402,6 @@ int runTestHarnessWithCheck( int argc, const char *argv[], int testNum, test_def
             gHasLong = 0;
     }
 
-    if( getenv( "OPENCL_1_0_DEVICE" ) )
-    {
-        char c_version[1024];
-        gIsOpenCL_1_0_Device = 1;
-        memset( c_version, 0, sizeof( c_version ) );
-
-        if( (err = clGetDeviceInfo( device, CL_DEVICE_OPENCL_C_VERSION, sizeof(c_version), c_version, NULL )) )
-        {
-            log_error( "FAILURE: unable to get CL_DEVICE_OPENCL_C_VERSION on 1.0 device. (%d)\n", err );
-            return EXIT_FAILURE;
-        }
-
-        if( 0 == strncmp( c_version, "OpenCL C 1.0 ", strlen( "OpenCL C 1.0 " ) ) )
-        {
-            gIsOpenCL_C_1_0_Device = 1;
-            log_info( "Device is a OpenCL C 1.0 device\n" );
-        }
-        else
-            log_info( "Device is a OpenCL 1.0 device, but supports OpenCL C 1.1\n" );
-    }
-
     cl_uint device_address_bits = 0;
     if( (err = clGetDeviceInfo( device, CL_DEVICE_ADDRESS_BITS, sizeof( device_address_bits ), &device_address_bits, NULL ) ))
     {
@@ -662,6 +640,19 @@ int parseAndCallCommandLineTests( int argc, const char *argv[], cl_device_id dev
         }
     }
 
+    if (std::any_of(resultTestList, resultTestList + testNum,
+                    [](test_status result) {
+                        switch (result)
+                        {
+                            case TEST_PASS:
+                            case TEST_SKIP: return false;
+                            case TEST_FAIL: return true;
+                        };
+                    }))
+    {
+        ret = EXIT_FAILURE;
+    }
+
     free( selectedTestList );
     free( resultTestList );
 

test_common/harness/testHarness.h

@@ -147,7 +147,6 @@ extern int      gFlushDenormsToZero;    // This is set to 1 if the device does n
 extern int      gInfNanSupport;         // This is set to 1 if the device supports infinities and NaNs
 extern int        gIsEmbedded;            // This is set to 1 if the device is an embedded device
 extern int        gHasLong;               // This is set to 1 if the device suppots long and ulong types in OpenCL C.
-extern int      gIsOpenCL_C_1_0_Device; // This is set to 1 if the device supports only OpenCL C 1.0.
 extern bool gCoreILProgram;
 
 #if ! defined( __APPLE__ )

test_conformance/CMakeLists.txt

@@ -49,7 +49,9 @@ add_subdirectory( subgroups )
 add_subdirectory( workgroups )
 add_subdirectory( pipes )
 add_subdirectory( device_timer )
-add_subdirectory( clcpp )
+if(KHRONOS_OFFLINE_COMPILER)
+    add_subdirectory( clcpp )
+endif()
 add_subdirectory( spirv_new )
 add_subdirectory( spir )
 

test_conformance/api/CMakeLists.txt

@@ -23,6 +23,7 @@ set(${MODULE_NAME}_SOURCES
          test_kernel_arg_info_compatibility.cpp
          test_null_buffer_arg.cpp
          test_mem_object_info.cpp
+         test_min_image_formats.cpp
          test_queue.cpp
          test_queue_hint.cpp
          test_queue_properties.cpp
@@ -32,6 +33,7 @@ set(${MODULE_NAME}_SOURCES
          test_context_destructor_callback.cpp
          test_mem_object_properties_queries.cpp
          test_queue_properties_queries.cpp
+         test_pipe_properties_queries.cpp
 )
 
 include(../CMakeCommon.txt)

test_conformance/api/main.cpp

@@ -127,6 +127,7 @@ test_definition test_list[] = {
     ADD_TEST_VERSION(buffer_properties_queries, Version(3, 0)),
     ADD_TEST_VERSION(image_properties_queries, Version(3, 0)),
     ADD_TEST_VERSION(queue_properties_queries, Version(3, 0)),
+    ADD_TEST_VERSION(pipe_properties_queries, Version(3, 0)),
 
     ADD_TEST_VERSION(consistency_svm, Version(3, 0)),
     ADD_TEST_VERSION(consistency_memory_model, Version(3, 0)),
@@ -142,6 +143,8 @@ test_definition test_list[] = {
     ADD_TEST_VERSION(consistency_subgroups, Version(3, 0)),
     ADD_TEST_VERSION(consistency_prog_ctor_dtor, Version(3, 0)),
     ADD_TEST_VERSION(consistency_3d_image_writes, Version(3, 0)),
+
+    ADD_TEST(min_image_formats),
 };
 
 const int test_num = ARRAY_SIZE(test_list);

test_conformance/api/procs.h

@@ -137,6 +137,8 @@ extern int test_queue_properties_queries(cl_device_id deviceID,
                                          cl_context context,
                                          cl_command_queue queue,
                                          int num_elements);
+int test_pipe_properties_queries(cl_device_id deviceID, cl_context context,
+                                 cl_command_queue queue, int num_elements);
 
 extern int test_consistency_svm(cl_device_id deviceID, cl_context context,
                                 cl_command_queue queue, int num_elements);
@@ -186,3 +188,6 @@ extern int test_consistency_3d_image_writes(cl_device_id deviceID,
                                             cl_context context,
                                             cl_command_queue queue,
                                             int num_elements);
+
+extern int test_min_image_formats(cl_device_id deviceID, cl_context context,
+                                  cl_command_queue queue, int num_elements);

test_conformance/api/test_api_min_max.cpp

@@ -136,7 +136,8 @@ int test_min_max_thread_dimensions(cl_device_id deviceID, cl_context context, cl
     }
 
     /* Create some I/O streams */
-    streams[0] = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_int) * 100, NULL, &error );
+    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_int) * 100, NULL, &error);
     if( streams[0] == NULL )
     {
         log_error("ERROR: Creating test array failed!\n");
@@ -321,7 +322,8 @@ int test_min_max_read_image_args(cl_device_id deviceID, cl_context context, cl_c
     test_error( error, "Failed to create the program and kernel.");
     free( programSrc );
 
-    result = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_float), NULL, &error);
+    result = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_float), NULL,
+                            &error);
     test_error( error, "clCreateBufer failed");
 
     /* Create some I/O streams */
@@ -692,7 +694,8 @@ int test_min_max_image_3d_width(cl_device_id deviceID, cl_context context, cl_co
     PASSIVE_REQUIRE_3D_IMAGE_SUPPORT( deviceID )
 
     /* Just get any ol format to test with */
-    error = get_8_bit_image_format( context, CL_MEM_OBJECT_IMAGE3D, CL_MEM_READ_WRITE, 0, &image_format_desc );
+    error = get_8_bit_image_format(context, CL_MEM_OBJECT_IMAGE3D,
+                                   CL_MEM_READ_ONLY, 0, &image_format_desc);
     test_error( error, "Unable to obtain suitable image format to test with!" );
 
     /* Get the max 2d image width */
@@ -748,7 +751,8 @@ int test_min_max_image_3d_height(cl_device_id deviceID, cl_context context, cl_c
     PASSIVE_REQUIRE_3D_IMAGE_SUPPORT( deviceID )
 
     /* Just get any ol format to test with */
-    error = get_8_bit_image_format( context, CL_MEM_OBJECT_IMAGE3D, CL_MEM_READ_WRITE, 0, &image_format_desc );
+    error = get_8_bit_image_format(context, CL_MEM_OBJECT_IMAGE3D,
+                                   CL_MEM_READ_ONLY, 0, &image_format_desc);
     test_error( error, "Unable to obtain suitable image format to test with!" );
 
     /* Get the max 2d image width */
@@ -805,7 +809,8 @@ int test_min_max_image_3d_depth(cl_device_id deviceID, cl_context context, cl_co
     PASSIVE_REQUIRE_3D_IMAGE_SUPPORT( deviceID )
 
     /* Just get any ol format to test with */
-    error = get_8_bit_image_format( context, CL_MEM_OBJECT_IMAGE3D, CL_MEM_READ_WRITE, 0, &image_format_desc );
+    error = get_8_bit_image_format(context, CL_MEM_OBJECT_IMAGE3D,
+                                   CL_MEM_READ_ONLY, 0, &image_format_desc);
     test_error( error, "Unable to obtain suitable image format to test with!" );
 
     /* Get the max 2d image width */
@@ -991,6 +996,7 @@ int test_min_max_parameter_size(cl_device_id deviceID, cl_context context, cl_co
     size_t decrement;
     cl_event event;
     cl_int event_status;
+    bool embeddedNoLong = gIsEmbedded && !gHasLong;
 
 
     /* Get the max param size */
@@ -1004,8 +1010,9 @@ int test_min_max_parameter_size(cl_device_id deviceID, cl_context context, cl_co
         return -1;
     }
 
-    /* The embedded profile does not require longs, so use ints */
+    /* The embedded profile without cles_khr_int64 extension does not require
-    if(gIsEmbedded)
+     * longs, so use ints */
+    if (embeddedNoLong)
         numberOfIntParametersToTry = numberExpected = (maxSize-sizeof(cl_mem))/sizeof(cl_int);
     else
         numberOfIntParametersToTry = numberExpected = (maxSize-sizeof(cl_mem))/sizeof(cl_long);
@@ -1021,7 +1028,7 @@ int test_min_max_parameter_size(cl_device_id deviceID, cl_context context, cl_co
         clMemWrapper mem;
         clKernelWrapper kernel;
 
-        if(gIsEmbedded)
+        if (embeddedNoLong)
         {
             log_info("Trying a kernel with %ld int arguments (%ld bytes) and one cl_mem (%ld bytes) for %ld bytes total.\n",
                      numberOfIntParametersToTry, sizeof(cl_int)*numberOfIntParametersToTry, sizeof(cl_mem),
@@ -1092,7 +1099,8 @@ int test_min_max_parameter_size(cl_device_id deviceID, cl_context context, cl_co
         /* Try to set a large argument to the kernel */
         retVal = 0;
 
-        mem = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_long), NULL, &error);
+        mem = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_long), NULL,
+                             &error);
         test_error(error, "clCreateBuffer failed");
 
         for (i=0; i<(int)numberOfIntParametersToTry; i++) {
@@ -1246,7 +1254,8 @@ int test_min_max_samplers(cl_device_id deviceID, cl_context context, cl_command_
     clMemWrapper image = create_image_2d( context, CL_MEM_READ_WRITE, &format, 16, 16, 0, NULL, &error );
     test_error( error, "Unable to create a test image" );
 
-    clMemWrapper stream = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE), 16, NULL, &error );
+    clMemWrapper stream =
+        clCreateBuffer(context, CL_MEM_READ_WRITE, 16, NULL, &error);
     test_error( error, "Unable to create test buffer" );
 
     error = clSetKernelArg( kernel, 0, sizeof( cl_mem ), &image );
@@ -1347,9 +1356,11 @@ int test_min_max_constant_buffer_size(cl_device_id deviceID, cl_context context,
             constantData[i] = (int)genrand_int32(d);
 
         clMemWrapper streams[3];
-        streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_COPY_HOST_PTR), sizeToAllocate, constantData, &error);
+        streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
+                                    sizeToAllocate, constantData, &error);
         test_error( error, "Creating test array failed" );
-        streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeToAllocate, NULL, &error);
+        streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeToAllocate,
+                                    NULL, &error);
         test_error( error, "Creating test array failed" );
 
 
@@ -1513,7 +1524,8 @@ int test_min_max_constant_args(cl_device_id deviceID, cl_context context, cl_com
     streams = new clMemWrapper[ maxArgs + 1 ];
     for( i = 0; i < maxArgs + 1; i++ )
     {
-        streams[i] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), individualBufferSize, NULL, &error);
+        streams[i] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                    individualBufferSize, NULL, &error);
         test_error( error, "Creating test array failed" );
     }
 
@@ -1721,9 +1733,11 @@ int test_min_max_local_mem_size(cl_device_id deviceID, cl_context context, cl_co
         localData[i] = (int)genrand_int32(d);
     free_mtdata(d); d = NULL;
 
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_COPY_HOST_PTR), sizeToAllocate, localData, &error);
+    streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, sizeToAllocate,
+                                localData, &error);
     test_error( error, "Creating test array failed" );
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeToAllocate, NULL, &error);
+    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeToAllocate,
+                                NULL, &error);
     test_error( error, "Creating test array failed" );
 
 

test_conformance/api/test_create_context_from_type.cpp

@@ -80,9 +80,11 @@ int test_create_context_from_type(cl_device_id deviceID, cl_context context, cl_
     }
 
     /* Create some I/O streams */
-    streams[0] = clCreateBuffer(context_to_test, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_float) * 10, NULL, &error);
+    streams[0] = clCreateBuffer(context_to_test, CL_MEM_READ_WRITE,
+                                sizeof(cl_float) * 10, NULL, &error);
     test_error( error, "Creating test array failed" );
-    streams[1] = clCreateBuffer(context_to_test, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_int) * 10, NULL, &error);
+    streams[1] = clCreateBuffer(context_to_test, CL_MEM_READ_WRITE,
+                                sizeof(cl_int) * 10, NULL, &error);
     test_error( error, "Creating test array failed" );
 
     /* Write some test data */

test_conformance/api/test_create_kernels.cpp

@@ -456,7 +456,8 @@ int test_enqueue_task(cl_device_id deviceID, cl_context context, cl_command_queu
 
     // Create args
     count = 100;
-    output = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof( cl_int ) * count, NULL, &error );
+    output = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int) * count,
+                            NULL, &error);
     test_error( error, "Unable to create output buffer" );
 
     error = clSetKernelArg( kernel, 0, sizeof( cl_mem ), &output );

test_conformance/api/test_kernel_arg_changes.cpp

@@ -74,14 +74,16 @@ int test_kernel_arg_changes(cl_device_id device, cl_context context, cl_command_
         sizes[ i ][ 0 ] = genrand_int32(seed) % (maxWidth/32) + 1;
         sizes[ i ][ 1 ] = genrand_int32(seed) % (maxHeight/32) + 1;
 
-        images[ i ] = create_image_2d( context, (cl_mem_flags)(CL_MEM_READ_ONLY),
+        images[i] = create_image_2d(context, CL_MEM_READ_ONLY, &imageFormat,
-                                     &imageFormat, sizes[ i ][ 0], sizes[ i ][ 1 ], 0, NULL, &error );
+                                    sizes[i][0], sizes[i][1], 0, NULL, &error);
         if( images[i] == NULL )
         {
             log_error("Failed to create image %d of size %d x %d (%s).\n", i, (int)sizes[i][0], (int)sizes[i][1], IGetErrorString( error ));
             return -1;
         }
-        results[ i ] = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof( cl_int ) * threads[0] * 2, NULL, &error );
+        results[i] =
+            clCreateBuffer(context, CL_MEM_READ_WRITE,
+                           sizeof(cl_int) * threads[0] * 2, NULL, &error);
         if( results[i] == NULL)
         {
             log_error("Failed to create array %d of size %d.\n", i, (int)threads[0]*2);

test_conformance/api/test_kernel_arg_multi_setup.cpp

@@ -66,24 +66,39 @@ int test_multi_arg_set(cl_device_id device, cl_context context, cl_command_queue
 
     // Create input streams
     initData[ 0 ] = create_random_data( vec1Type, d, (unsigned int)threads[ 0 ] * vec1Size );
-    streams[ 0 ] = clCreateBuffer( context, (cl_mem_flags)( CL_MEM_COPY_HOST_PTR ), get_explicit_type_size( vec1Type ) * threads[0] * vec1Size, initData[ 0 ], &error );
+    streams[0] =
+        clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
+                       get_explicit_type_size(vec1Type) * threads[0] * vec1Size,
+                       initData[0], &error);
     test_error( error, "Unable to create testing stream" );
 
     initData[ 1 ] = create_random_data( vec2Type, d, (unsigned int)threads[ 0 ] * vec2Size );
-    streams[ 1 ] = clCreateBuffer( context, (cl_mem_flags)( CL_MEM_COPY_HOST_PTR ), get_explicit_type_size( vec2Type ) * threads[0] * vec2Size, initData[ 1 ], &error );
+    streams[1] =
+        clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
+                       get_explicit_type_size(vec2Type) * threads[0] * vec2Size,
+                       initData[1], &error);
     test_error( error, "Unable to create testing stream" );
 
     initData[ 2 ] = create_random_data( vec3Type, d, (unsigned int)threads[ 0 ] * vec3Size );
-    streams[ 2 ] = clCreateBuffer( context, (cl_mem_flags)( CL_MEM_COPY_HOST_PTR ), get_explicit_type_size( vec3Type ) * threads[0] * vec3Size, initData[ 2 ], &error );
+    streams[2] =
+        clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
+                       get_explicit_type_size(vec3Type) * threads[0] * vec3Size,
+                       initData[2], &error);
     test_error( error, "Unable to create testing stream" );
 
-    streams[ 3 ] = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE),  get_explicit_type_size( vec1Type ) * threads[0] * vec1Size, NULL, &error );
+    streams[3] = clCreateBuffer(
+        context, CL_MEM_READ_WRITE,
+        get_explicit_type_size(vec1Type) * threads[0] * vec1Size, NULL, &error);
     test_error( error, "Unable to create testing stream" );
 
-    streams[ 4 ] = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE),  get_explicit_type_size( vec2Type ) * threads[0] * vec2Size, NULL, &error );
+    streams[4] = clCreateBuffer(
+        context, CL_MEM_READ_WRITE,
+        get_explicit_type_size(vec2Type) * threads[0] * vec2Size, NULL, &error);
     test_error( error, "Unable to create testing stream" );
 
-    streams[ 5 ] = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE),  get_explicit_type_size( vec3Type ) * threads[0] * vec3Size, NULL, &error );
+    streams[5] = clCreateBuffer(
+        context, CL_MEM_READ_WRITE,
+        get_explicit_type_size(vec3Type) * threads[0] * vec3Size, NULL, &error);
     test_error( error, "Unable to create testing stream" );
 
     // Set the arguments

test_conformance/api/test_kernels.cpp

@@ -192,10 +192,10 @@ int test_execute_kernel_local_sizes(cl_device_id deviceID, cl_context context, c
     }
 
     /* Create some I/O streams */
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),
+    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
                                 sizeof(cl_float) * num_elements, NULL, &error);
     test_error( error, "Creating test array failed" );
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),
+    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
                                 sizeof(cl_int) * num_elements, NULL, &error);
     test_error( error, "Creating test array failed" );
 
@@ -323,10 +323,10 @@ int test_set_kernel_arg_by_index(cl_device_id deviceID, cl_context context, cl_c
     }
 
     /* Create some I/O streams */
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),
+    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
                                 sizeof(cl_float) * num_elements, NULL, &error);
     test_error( error, "Creating test array failed" );
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),
+    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
                                 sizeof(cl_int) * num_elements, NULL, &error);
     test_error( error, "Creating test array failed" );
 
@@ -412,15 +412,15 @@ int test_set_kernel_arg_constant(cl_device_id deviceID, cl_context context, cl_c
     }
     free_mtdata(d); d = NULL;
 
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_COPY_HOST_PTR),
+    streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
                                 sizeof(cl_int) * num_elements,
                                 randomTestDataA.data(), &error);
     test_error( error, "Creating test array failed" );
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_COPY_HOST_PTR),
+    streams[1] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
                                 sizeof(cl_int) * num_elements,
                                 randomTestDataB.data(), &error);
     test_error( error, "Creating test array failed" );
-    streams[2] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),
+    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE,
                                 sizeof(cl_int) * num_elements, NULL, &error);
     test_error( error, "Creating test array failed" );
 
@@ -496,11 +496,11 @@ int test_set_kernel_arg_struct_array(cl_device_id deviceID, cl_context context,
     }
     free_mtdata(d); d = NULL;
 
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_COPY_HOST_PTR),
+    streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
                                 sizeof(image_pair_t) * num_elements,
                                 (void *)image_pair.data(), &error);
     test_error( error, "Creating test array failed" );
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),
+    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
                                 sizeof(cl_int) * num_elements, NULL, &error);
     test_error( error, "Creating test array failed" );
 
@@ -593,11 +593,11 @@ int test_kernel_global_constant(cl_device_id deviceID, cl_context context, cl_co
     }
     free_mtdata(d); d = NULL;
 
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_COPY_HOST_PTR),
+    streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
                                 sizeof(cl_int) * num_elements,
                                 randomTestDataA.data(), &error);
     test_error( error, "Creating test array failed" );
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),
+    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
                                 sizeof(cl_int) * num_elements, NULL, &error);
     test_error( error, "Creating test array failed" );
 

test_conformance/api/test_mem_object_properties_queries.cpp

@@ -22,16 +22,18 @@
 typedef enum
 {
     image,
-    buffer
+    image_with_properties,
-} mem_obj_type;
+    buffer,
+    buffer_with_properties,
+    subbuffer,
+} test_type;
 
 struct test_data
 {
-    mem_obj_type obj_t;
+    test_type type;
     std::vector<cl_mem_properties> properties;
     std::string description;
-    std::string src;
+    cl_kernel kernel;
-    std::string kernel_name;
 };
 
 static int create_object_and_check_properties(cl_context context,
@@ -43,47 +45,94 @@ static int create_object_and_check_properties(cl_context context,
 {
     cl_int error = CL_SUCCESS;
 
-    if (test_case.obj_t == image)
+    switch (test_case.type)
     {
-        cl_image_format format;
+        case image: {
-        format.image_channel_order = CL_RGBA;
+            cl_image_format format = { 0 };
-        format.image_channel_data_type = CL_UNSIGNED_INT32;
+            format.image_channel_order = CL_RGBA;
-        cl_image_desc desc;
+            format.image_channel_data_type = CL_UNSIGNED_INT32;
-        memset(&desc, 0x0, sizeof(cl_image_desc));
+            test_object = clCreateImage2D(context, flags, &format, size_x,
-        desc.image_type = CL_MEM_OBJECT_IMAGE2D;
+                                          size_y, 0, local_data.data(), &error);
-        desc.image_width = size_x;
+            test_error(error, "clCreateImage2D failed");
-        desc.image_height = size_y;
+        }
+        break;
+        case image_with_properties: {
+            cl_image_format format = { 0 };
+            format.image_channel_order = CL_RGBA;
+            format.image_channel_data_type = CL_UNSIGNED_INT32;
+            cl_image_desc desc = { 0 };
+            desc.image_type = CL_MEM_OBJECT_IMAGE2D;
+            desc.image_width = size_x;
+            desc.image_height = size_y;
 
-        if (test_case.properties.size() == 0)
+            if (test_case.properties.size() == 0)
-        {
+            {
-            test_object =
+                test_object = clCreateImageWithProperties(
-                clCreateImageWithProperties(context, NULL, flags, &format,
+                    context, NULL, flags, &format, &desc, local_data.data(),
-                                            &desc, local_data.data(), &error);
+                    &error);
+            }
+            else
+            {
+                test_object = clCreateImageWithProperties(
+                    context, test_case.properties.data(), flags, &format, &desc,
+                    local_data.data(), &error);
+            }
+            test_error(error, "clCreateImageWithProperties failed");
         }
-        else
+        break;
-        {
+        case buffer: {
-            test_object = clCreateImageWithProperties(
+            test_object = clCreateBuffer(context, flags,
-                context, test_case.properties.data(), flags, &format, &desc,
+                                         local_data.size() * sizeof(cl_uint),
-                local_data.data(), &error);
+                                         local_data.data(), &error);
+            test_error(error, "clCreateBuffer failed");
         }
-        test_error(error, "clCreateImageWithProperties failed");
+        case buffer_with_properties: {
-    }
+            if (test_case.properties.size() == 0)
-    if (test_case.obj_t == buffer)
+            {
-    {
+                test_object = clCreateBufferWithProperties(
-        if (test_case.properties.size() == 0)
+                    context, NULL, flags, local_data.size() * sizeof(cl_uint),
-        {
+                    local_data.data(), &error);
-            test_object = clCreateBufferWithProperties(
+            }
-                context, NULL, flags, local_data.size() * sizeof(cl_uint),
+            else
-                local_data.data(), &error);
+            {
-        }
+                test_object = clCreateBufferWithProperties(
-        else
+                    context, test_case.properties.data(), flags,
-        {
+                    local_data.size() * sizeof(cl_uint), local_data.data(),
-            test_object = clCreateBufferWithProperties(
+                    &error);
-                context, test_case.properties.data(), flags,
+            }
-                local_data.size() * sizeof(cl_uint), local_data.data(), &error);
+            test_error(error, "clCreateBufferWithProperties failed.");
         }
+        break;
+        case subbuffer: {
+            clMemWrapper parent_object;
+            if (test_case.properties.size() == 0)
+            {
+                parent_object = clCreateBufferWithProperties(
+                    context, NULL, flags, local_data.size() * sizeof(cl_uint),
+                    local_data.data(), &error);
+            }
+            else
+            {
+                parent_object = clCreateBufferWithProperties(
+                    context, test_case.properties.data(), flags,
+                    local_data.size() * sizeof(cl_uint), local_data.data(),
+                    &error);
+            }
+            test_error(error, "clCreateBufferWithProperties failed.");
 
-        test_error(error, "clCreateBufferWithProperties failed.");
+            cl_mem_flags subbuffer_flags = flags
+                & (CL_MEM_READ_WRITE | CL_MEM_READ_ONLY | CL_MEM_WRITE_ONLY);
+
+            cl_buffer_region region = { 0 };
+            region.origin = 0;
+            region.size = local_data.size() * sizeof(cl_uint);
+            test_object = clCreateSubBuffer(parent_object, subbuffer_flags,
+                                            CL_BUFFER_CREATE_TYPE_REGION,
+                                            &region, &error);
+            test_error(error, "clCreateSubBuffer failed.");
+        }
+        break;
+        default: log_error("Unknown test type!"); return TEST_FAIL;
     }
 
     std::vector<cl_mem_properties> check_properties;
@@ -94,6 +143,22 @@ static int create_object_and_check_properties(cl_context context,
     test_error(error,
                "clGetMemObjectInfo failed asking for CL_MEM_PROPERTIES size.");
 
+    // Buffers, subbuffers, and images must return no properties.
+    if (test_case.type == buffer || test_case.type == subbuffer
+        || test_case.type == image)
+    {
+        if (set_size == 0)
+        {
+            return TEST_PASS;
+        }
+        else
+        {
+            log_error("Queried properties must have size equal to zero for "
+                      "buffers, subbuffers, and images.");
+            return TEST_FAIL;
+        }
+    }
+
     if (set_size == 0 && test_case.properties.size() == 0)
     {
         return TEST_PASS;
@@ -123,8 +188,6 @@ static int run_test_query_properties(cl_context context, cl_command_queue queue,
     int error = CL_SUCCESS;
     log_info("\nTC description: %s\n", test_case.description.c_str());
 
-    clProgramWrapper program;
-    clKernelWrapper kernel;
     clMemWrapper obj_src;
     clMemWrapper obj_dst;
     clEventWrapper event;
@@ -144,54 +207,55 @@ static int run_test_query_properties(cl_context context, cl_command_queue queue,
     generate_random_data(kUInt, size, init_generator, dst_data.data());
     free_mtdata(init_generator);
     init_generator = NULL;
-    const char* kernel_src = test_case.src.c_str();
-    error =
-        create_single_kernel_helper(context, &program, &kernel, 1, &kernel_src,
-                                    test_case.kernel_name.c_str());
 
-    test_error(error, "create_single_kernel_helper failed");
+    flags = CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR;
-
-    flags = (cl_mem_flags)(CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR);
     error = create_object_and_check_properties(context, obj_src, test_case,
                                                flags, src_data, size_x, size_y);
     test_error(error, "create_object_and_check_properties obj_src failed.");
 
-    flags = (cl_mem_flags)(CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR);
+    flags = CL_MEM_WRITE_ONLY | CL_MEM_COPY_HOST_PTR;
     error = create_object_and_check_properties(context, obj_dst, test_case,
                                                flags, dst_data, size_x, size_y);
     test_error(error, "create_object_and_check_properties obj_dst failed.");
 
-    error = clSetKernelArg(kernel, 0, sizeof(obj_src), &obj_src);
+    error = clSetKernelArg(test_case.kernel, 0, sizeof(obj_src), &obj_src);
     test_error(error, "clSetKernelArg 0 failed.");
 
-    error = clSetKernelArg(kernel, 1, sizeof(obj_dst), &obj_dst);
+    error = clSetKernelArg(test_case.kernel, 1, sizeof(obj_dst), &obj_dst);
     test_error(error, "clSetKernelArg 1 failed.");
-    if (test_case.obj_t == image)
+    switch (test_case.type)
     {
-        error = clEnqueueNDRangeKernel(queue, kernel, 2, NULL, global_dim, NULL,
+        case image:
-                                       0, NULL, &event);
+        case image_with_properties: {
-        test_error(error, "clEnqueueNDRangeKernel failed.");
+            error = clEnqueueNDRangeKernel(queue, test_case.kernel, 2, NULL,
+                                           global_dim, NULL, 0, NULL, &event);
+            test_error(error, "clEnqueueNDRangeKernel failed.");
 
-        error = clWaitForEvents(1, &event);
+            error = clWaitForEvents(1, &event);
-        test_error(error, "clWaitForEvents failed.");
+            test_error(error, "clWaitForEvents failed.");
 
-        error = clEnqueueReadImage(queue, obj_dst, CL_TRUE, origin, region, 0,
+            error = clEnqueueReadImage(queue, obj_dst, CL_TRUE, origin, region,
-                                   0, dst_data.data(), 0, NULL, NULL);
+                                       0, 0, dst_data.data(), 0, NULL, NULL);
-        test_error(error, "clEnqueueReadImage failed.");
+            test_error(error, "clEnqueueReadImage failed.");
-    }
+        }
-    if (test_case.obj_t == buffer)
+        break;
-    {
+        case buffer:
-        error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &size, NULL, 0,
+        case buffer_with_properties:
-                                       NULL, &event);
+        case subbuffer: {
-        test_error(error, "clEnqueueNDRangeKernel failed.");
+            error = clEnqueueNDRangeKernel(queue, test_case.kernel, 1, NULL,
+                                           &size, NULL, 0, NULL, &event);
+            test_error(error, "clEnqueueNDRangeKernel failed.");
 
-        error = clWaitForEvents(1, &event);
+            error = clWaitForEvents(1, &event);
-        test_error(error, "clWaitForEvents failed.");
+            test_error(error, "clWaitForEvents failed.");
 
-        error = clEnqueueReadBuffer(queue, obj_dst, CL_TRUE, 0,
+            error = clEnqueueReadBuffer(queue, obj_dst, CL_TRUE, 0,
-                                    dst_data.size() * sizeof(cl_uint),
+                                        dst_data.size() * sizeof(cl_uint),
-                                    dst_data.data(), 0, NULL, NULL);
+                                        dst_data.data(), 0, NULL, NULL);
-        test_error(error, "clEnqueueReadBuffer failed.");
+            test_error(error, "clEnqueueReadBuffer failed.");
+        }
+        break;
+        default: log_error("Unknown test type!"); return TEST_FAIL;
     }
 
     for (size_t i = 0; i < size; ++i)
@@ -223,21 +287,31 @@ int test_image_properties_queries(cl_device_id deviceID, cl_context context,
         return TEST_SKIPPED_ITSELF;
     }
 
+    clProgramWrapper program;
+    clKernelWrapper kernel;
+
+    const char* kernel_src = R"CLC(
+        __kernel void data_copy(read_only image2d_t src, write_only image2d_t dst)
+        {
+            int tid_x = get_global_id(0);
+            int tid_y = get_global_id(1);
+            int2 coords = (int2)(tid_x, tid_y);
+            uint4 val = read_imageui(src, coords);
+            write_imageui(dst, coords, val);
+
+        }
+        )CLC";
+
+    error = create_single_kernel_helper(context, &program, &kernel, 1,
+                                        &kernel_src, "data_copy");
+    test_error(error, "create_single_kernel_helper failed");
+
     std::vector<test_data> test_cases;
-    std::string test_kernel = { "__kernel void data_copy(read_only image2d_t "
+    test_cases.push_back({ image, {}, "regular image", kernel });
-                                "src, write_only image2d_t dst)\n"
-                                "{\n"
-                                "    int tid_x = get_global_id(0);\n"
-                                "    int tid_y = get_global_id(1);\n"
-                                "    int2 coords = (int2)(tid_x, tid_y);\n"
-                                "    uint4 val = read_imageui(src, coords);\n"
-                                "    write_imageui(dst, coords, val);\n"
-                                "\n"
-                                "}\n" };
     test_cases.push_back(
-        { image, { 0 }, "image, 0 properties", test_kernel, "data_copy" });
+        { image_with_properties, { 0 }, "image, 0 properties", kernel });
     test_cases.push_back(
-        { image, {}, "image, NULL properties", test_kernel, "data_copy" });
+        { image_with_properties, {}, "image, NULL properties", kernel });
 
     for (auto test_case : test_cases)
     {
@@ -251,20 +325,33 @@ int test_buffer_properties_queries(cl_device_id deviceID, cl_context context,
                                    cl_command_queue queue, int num_elements)
 {
     int error = CL_SUCCESS;
+
+    clProgramWrapper program;
+    clKernelWrapper kernel;
+
+    const char* kernel_src = R"CLC(
+        __kernel void data_copy(__global int *src, __global int *dst)
+        {
+            int  tid = get_global_id(0);
+
+            dst[tid] = src[tid];
+
+        }
+        )CLC";
+    error = create_single_kernel_helper(context, &program, &kernel, 1,
+                                        &kernel_src, "data_copy");
+    test_error(error, "create_single_kernel_helper failed");
+
     std::vector<test_data> test_cases;
-    std::string test_kernel = {
+    test_cases.push_back({ buffer, {}, "regular buffer", kernel });
-        "__kernel void data_copy(__global int *src, __global int *dst)\n"
-        "{\n"
-        "    int  tid = get_global_id(0);\n"
-        "\n"
-        "    dst[tid] = src[tid];\n"
-        "\n"
-        "}\n"
-    };
     test_cases.push_back(
-        { buffer, { 0 }, "buffer, 0 properties", test_kernel, "data_copy" });
+        { buffer_with_properties, { 0 }, "buffer with 0 properties", kernel });
     test_cases.push_back(
-        { buffer, {}, "buffer, NULL properties", test_kernel, "data_copy" });
+        { buffer_with_properties, {}, "buffer with NULL properties", kernel });
+    test_cases.push_back(
+        { subbuffer, { 0 }, "subbuffer with 0 properties", kernel });
+    test_cases.push_back(
+        { subbuffer, {}, "subbuffer with NULL properties", kernel });
 
     for (auto test_case : test_cases)
     {

test_conformance/api/test_min_image_formats.cpp

@@ -0,0 +1,133 @@
+//
+// Copyright (c) 2017 The Khronos Group Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+#include "testBase.h"
+
+int test_min_image_formats(cl_device_id device, cl_context context,
+                           cl_command_queue queue, int num_elements)
+{
+    int missingFormats = 0;
+
+    cl_int error = CL_SUCCESS;
+
+    Version version = get_device_cl_version(device);
+
+    cl_bool supports_images = CL_FALSE;
+    error = clGetDeviceInfo(device, CL_DEVICE_IMAGE_SUPPORT,
+                            sizeof(supports_images), &supports_images, NULL);
+    test_error(error, "clGetDeviceInfo for CL_DEVICE_IMAGE_SUPPORT failed");
+
+    if (supports_images == CL_FALSE)
+    {
+        log_info("No image support on current device - skipped\n");
+        return TEST_SKIPPED_ITSELF;
+    }
+
+    const cl_mem_object_type image_types[] = {
+        CL_MEM_OBJECT_IMAGE1D,       CL_MEM_OBJECT_IMAGE1D_BUFFER,
+        CL_MEM_OBJECT_IMAGE2D,       CL_MEM_OBJECT_IMAGE3D,
+        CL_MEM_OBJECT_IMAGE1D_ARRAY, CL_MEM_OBJECT_IMAGE2D_ARRAY,
+    };
+    const cl_mem_flags mem_flags[] = {
+        CL_MEM_READ_ONLY,
+        CL_MEM_WRITE_ONLY,
+        CL_MEM_KERNEL_READ_AND_WRITE,
+    };
+
+    cl_bool supports_read_write_images = CL_FALSE;
+    if (version >= Version(3, 0))
+    {
+        cl_uint maxReadWriteImageArgs = 0;
+        error = clGetDeviceInfo(device, CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS,
+                                sizeof(maxReadWriteImageArgs),
+                                &maxReadWriteImageArgs, NULL);
+        test_error(error,
+                   "Unable to query "
+                   "CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS");
+
+        // read-write images are supported if MAX_READ_WRITE_IMAGE_ARGS is
+        // nonzero
+        supports_read_write_images =
+            maxReadWriteImageArgs != 0 ? CL_TRUE : CL_FALSE;
+    }
+    else if (version >= Version(2, 0))
+    {
+        // read-write images are required for OpenCL 2.x
+        supports_read_write_images = CL_TRUE;
+    }
+
+    int supports_3D_image_writes =
+        is_extension_available(device, "cl_khr_3d_image_writes");
+
+    for (int t = 0; t < ARRAY_SIZE(image_types); t++)
+    {
+        const cl_mem_object_type type = image_types[t];
+        log_info("    testing %s...\n", convert_image_type_to_string(type));
+        for (int f = 0; f < ARRAY_SIZE(mem_flags); f++)
+        {
+            const cl_mem_flags flags = mem_flags[f];
+            const char* testTypeString = flags == CL_MEM_READ_ONLY
+                ? "read-only"
+                : flags == CL_MEM_WRITE_ONLY
+                    ? "write only"
+                    : flags == CL_MEM_KERNEL_READ_AND_WRITE ? "read and write"
+                                                            : "unknown???";
+
+            if (flags == CL_MEM_KERNEL_READ_AND_WRITE
+                && !supports_read_write_images)
+            {
+                continue;
+            }
+
+            if (type == CL_MEM_OBJECT_IMAGE3D && flags != CL_MEM_READ_ONLY
+                && !supports_3D_image_writes)
+            {
+                continue;
+            }
+
+            cl_uint numImageFormats = 0;
+            error = clGetSupportedImageFormats(context, flags, type, 0, NULL,
+                                               &numImageFormats);
+            test_error(error, "Unable to query number of image formats");
+
+            std::vector<cl_image_format> supportedFormats(numImageFormats);
+            if (numImageFormats != 0)
+            {
+                error = clGetSupportedImageFormats(
+                    context, flags, type, supportedFormats.size(),
+                    supportedFormats.data(), NULL);
+                test_error(error, "Unable to query image formats");
+            }
+
+            std::vector<cl_image_format> requiredFormats;
+            build_required_image_formats(flags, type, device, requiredFormats);
+
+            for (auto& format : requiredFormats)
+            {
+                if (!find_format(supportedFormats.data(),
+                                 supportedFormats.size(), &format))
+                {
+                    log_error(
+                        "Missing required %s format %s + %s.\n", testTypeString,
+                        GetChannelOrderName(format.image_channel_order),
+                        GetChannelTypeName(format.image_channel_data_type));
+                    ++missingFormats;
+                }
+            }
+        }
+    }
+
+    return missingFormats == 0 ? TEST_PASS : TEST_FAIL;
+}

test_conformance/api/test_pipe_properties_queries.cpp

@@ -0,0 +1,100 @@
+//
+// Copyright (c) 2020 The Khronos Group Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+#include "testBase.h"
+#include "harness/typeWrappers.h"
+
+#include <vector>
+
+struct test_query_pipe_properties_data
+{
+    std::vector<cl_pipe_properties> properties;
+    std::string description;
+};
+
+static int create_pipe_and_check_array_properties(
+    cl_context context, const test_query_pipe_properties_data& test_case)
+{
+    log_info("TC description: %s\n", test_case.description.c_str());
+
+    cl_int error = CL_SUCCESS;
+
+    clMemWrapper test_pipe;
+
+    if (test_case.properties.size() > 0)
+    {
+        test_pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, 4, 4,
+                                 test_case.properties.data(), &error);
+        test_error(error, "clCreatePipe failed");
+    }
+    else
+    {
+        test_pipe =
+            clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, 4, 4, NULL, &error);
+        test_error(error, "clCreatePipe failed");
+    }
+
+    std::vector<cl_pipe_properties> check_properties;
+    size_t set_size = 0;
+
+    error = clGetPipeInfo(test_pipe, CL_PIPE_PROPERTIES, 0, NULL, &set_size);
+    test_error(error,
+               "clGetPipeInfo failed asking for "
+               "CL_PIPE_PROPERTIES size.");
+
+    if (set_size == 0 && test_case.properties.size() == 0)
+    {
+        return TEST_PASS;
+    }
+    if (set_size != test_case.properties.size() * sizeof(cl_pipe_properties))
+    {
+        log_error("ERROR: CL_PIPE_PROPERTIES size is %d, expected %d.\n",
+                  set_size,
+                  test_case.properties.size() * sizeof(cl_pipe_properties));
+        return TEST_FAIL;
+    }
+
+    log_error("Unexpected test case size.  This test needs to be updated to "
+              "compare pipe properties.\n");
+    return TEST_FAIL;
+}
+
+int test_pipe_properties_queries(cl_device_id deviceID, cl_context context,
+                                 cl_command_queue queue, int num_elements)
+{
+    cl_int error = CL_SUCCESS;
+
+    cl_bool pipeSupport = CL_FALSE;
+    error = clGetDeviceInfo(deviceID, CL_DEVICE_PIPE_SUPPORT,
+                            sizeof(pipeSupport), &pipeSupport, NULL);
+    test_error(error, "Unable to query CL_DEVICE_PIPE_SUPPORT");
+
+    if (pipeSupport == CL_FALSE)
+    {
+        return TEST_SKIPPED_ITSELF;
+    }
+
+    int result = TEST_PASS;
+
+    std::vector<test_query_pipe_properties_data> test_cases;
+    test_cases.push_back({ {}, "NULL properties" });
+
+    for (auto test_case : test_cases)
+    {
+        result |= create_pipe_and_check_array_properties(context, test_case);
+    }
+
+    return result;
+}

test_conformance/api/test_retain_program.cpp

@@ -68,9 +68,11 @@ int test_release_during_execute( cl_device_id deviceID, cl_context context, cl_c
         return -1;
     }
 
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_float) * 10, NULL, &error);
+    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_float) * 10, NULL, &error);
     test_error( error, "Creating test array failed" );
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_int) * 10, NULL, &error);
+    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int) * 10,
+                                NULL, &error);
     test_error( error, "Creating test array failed" );
 
     /* Set the arguments */

test_conformance/atomics/test_atomics.cpp

@@ -243,13 +243,17 @@ int test_atomic_function(cl_device_id deviceID, cl_context context, cl_command_q
     for( size_t i = 0; i < numDestItems; i++ )
         memcpy( destItems + i * typeSize, startValue, typeSize );
 
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_COPY_HOST_PTR), typeSize * numDestItems, destItems, NULL);
+    streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
+                                typeSize * numDestItems, destItems, NULL);
     if (!streams[0])
     {
         log_error("ERROR: Creating output array failed!\n");
         return -1;
     }
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(( startRefValues != NULL ? CL_MEM_COPY_HOST_PTR : CL_MEM_READ_WRITE )), typeSize * threadSize, startRefValues, NULL);
+    streams[1] = clCreateBuffer(
+        context,
+        ((startRefValues != NULL ? CL_MEM_COPY_HOST_PTR : CL_MEM_READ_WRITE)),
+        typeSize * threadSize, startRefValues, NULL);
     if (!streams[1])
     {
         log_error("ERROR: Creating reference array failed!\n");

test_conformance/atomics/test_indexed_cases.cpp

@@ -64,12 +64,12 @@ int test_atomic_add_index(cl_device_id deviceID, cl_context context, cl_command_
              (int)numGlobalThreads, (int)numLocalThreads);
 
     // Create the counter that will keep track of where each thread writes.
-    counter = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),
+    counter = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int) * 1,
-                                   sizeof(cl_int) * 1, NULL, NULL);
+                             NULL, NULL);
     // Create the counters that will hold the results of each thread writing
     // its ID into a (hopefully) unique location.
-    counters = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),
+    counters = clCreateBuffer(context, CL_MEM_READ_WRITE,
-                                    sizeof(cl_int) * numGlobalThreads, NULL, NULL);
+                              sizeof(cl_int) * numGlobalThreads, NULL, NULL);
 
     // Reset all those locations to -1 to indciate they have not been used.
     cl_int *values = (cl_int*) malloc(sizeof(cl_int)*numGlobalThreads);
@@ -175,12 +175,15 @@ int add_index_bin_test(size_t *global_threads, cl_command_queue queue, cl_contex
              (int)global_threads[0], (int)local_threads[0]);
 
     // Allocate our storage
-    cl_mem bin_counters = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),
+    cl_mem bin_counters =
-                                        sizeof(cl_int) * number_of_bins, NULL, NULL);
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
-    cl_mem bins = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),
+                       sizeof(cl_int) * number_of_bins, NULL, NULL);
-                                sizeof(cl_int) * number_of_bins*max_counts_per_bin, NULL, NULL);
+    cl_mem bins = clCreateBuffer(
-    cl_mem bin_assignments = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_ONLY),
+        context, CL_MEM_READ_WRITE,
-                                           sizeof(cl_int) * number_of_items, NULL, NULL);
+        sizeof(cl_int) * number_of_bins * max_counts_per_bin, NULL, NULL);
+    cl_mem bin_assignments =
+        clCreateBuffer(context, CL_MEM_READ_ONLY,
+                       sizeof(cl_int) * number_of_items, NULL, NULL);
 
     if (bin_counters == NULL) {
         log_error("add_index_bin_test FAILED to allocate bin_counters.\n");

test_conformance/basic/CMakeLists.txt

@@ -39,6 +39,7 @@ set(${MODULE_NAME}_SOURCES
     test_async_copy.cpp
     test_sizeof.cpp
     test_vector_creation.cpp
+    test_vector_swizzle.cpp
     test_vec_type_hint.cpp
     test_numeric_constants.cpp
     test_constant_source.cpp

test_conformance/basic/main.cpp

@@ -26,130 +26,131 @@
 #include "procs.h"
 
 test_definition test_list[] = {
-    ADD_TEST( hostptr ),
+    ADD_TEST(hostptr),
-    ADD_TEST( fpmath_float ),
+    ADD_TEST(fpmath_float),
-    ADD_TEST( fpmath_float2 ),
+    ADD_TEST(fpmath_float2),
-    ADD_TEST( fpmath_float4 ),
+    ADD_TEST(fpmath_float4),
-    ADD_TEST( intmath_int ),
+    ADD_TEST(intmath_int),
-    ADD_TEST( intmath_int2 ),
+    ADD_TEST(intmath_int2),
-    ADD_TEST( intmath_int4 ),
+    ADD_TEST(intmath_int4),
-    ADD_TEST( intmath_long ),
+    ADD_TEST(intmath_long),
-    ADD_TEST( intmath_long2 ),
+    ADD_TEST(intmath_long2),
-    ADD_TEST( intmath_long4 ),
+    ADD_TEST(intmath_long4),
-    ADD_TEST( hiloeo ),
+    ADD_TEST(hiloeo),
-    ADD_TEST( if ),
+    ADD_TEST(if),
-    ADD_TEST( sizeof ),
+    ADD_TEST(sizeof),
-    ADD_TEST( loop ),
+    ADD_TEST(loop),
-    ADD_TEST( pointer_cast ),
+    ADD_TEST(pointer_cast),
-    ADD_TEST( local_arg_def ),
+    ADD_TEST(local_arg_def),
-    ADD_TEST( local_kernel_def ),
+    ADD_TEST(local_kernel_def),
-    ADD_TEST( local_kernel_scope ),
+    ADD_TEST(local_kernel_scope),
-    ADD_TEST( constant ),
+    ADD_TEST(constant),
-    ADD_TEST( constant_source ),
+    ADD_TEST(constant_source),
-    ADD_TEST( readimage ),
+    ADD_TEST(readimage),
-    ADD_TEST( readimage_int16 ),
+    ADD_TEST(readimage_int16),
-    ADD_TEST( readimage_fp32 ),
+    ADD_TEST(readimage_fp32),
-    ADD_TEST( writeimage ),
+    ADD_TEST(writeimage),
-    ADD_TEST( writeimage_int16 ),
+    ADD_TEST(writeimage_int16),
-    ADD_TEST( writeimage_fp32 ),
+    ADD_TEST(writeimage_fp32),
-    ADD_TEST( mri_one ),
+    ADD_TEST(mri_one),
 
-    ADD_TEST( mri_multiple ),
+    ADD_TEST(mri_multiple),
-    ADD_TEST( image_r8 ),
+    ADD_TEST(image_r8),
-    ADD_TEST( barrier ),
+    ADD_TEST(barrier),
-    ADD_TEST_VERSION( wg_barrier, Version(2, 0) ),
+    ADD_TEST_VERSION(wg_barrier, Version(2, 0)),
-    ADD_TEST( int2float ),
+    ADD_TEST(int2float),
-    ADD_TEST( float2int ),
+    ADD_TEST(float2int),
-    ADD_TEST( imagereadwrite ),
+    ADD_TEST(imagereadwrite),
-    ADD_TEST( imagereadwrite3d ),
+    ADD_TEST(imagereadwrite3d),
-    ADD_TEST( readimage3d ),
+    ADD_TEST(readimage3d),
-    ADD_TEST( readimage3d_int16 ),
+    ADD_TEST(readimage3d_int16),
-    ADD_TEST( readimage3d_fp32 ),
+    ADD_TEST(readimage3d_fp32),
-    ADD_TEST( bufferreadwriterect ),
+    ADD_TEST(bufferreadwriterect),
-    ADD_TEST( arrayreadwrite ),
+    ADD_TEST(arrayreadwrite),
-    ADD_TEST( arraycopy ),
+    ADD_TEST(arraycopy),
-    ADD_TEST( imagearraycopy ),
+    ADD_TEST(imagearraycopy),
-    ADD_TEST( imagearraycopy3d ),
+    ADD_TEST(imagearraycopy3d),
-    ADD_TEST( imagecopy ),
+    ADD_TEST(imagecopy),
-    ADD_TEST( imagecopy3d ),
+    ADD_TEST(imagecopy3d),
-    ADD_TEST( imagerandomcopy ),
+    ADD_TEST(imagerandomcopy),
-    ADD_TEST( arrayimagecopy ),
+    ADD_TEST(arrayimagecopy),
-    ADD_TEST( arrayimagecopy3d ),
+    ADD_TEST(arrayimagecopy3d),
-    ADD_TEST( imagenpot ),
+    ADD_TEST(imagenpot),
 
-    ADD_TEST( vload_global ),
+    ADD_TEST(vload_global),
-    ADD_TEST( vload_local ),
+    ADD_TEST(vload_local),
-    ADD_TEST( vload_constant ),
+    ADD_TEST(vload_constant),
-    ADD_TEST( vload_private ),
+    ADD_TEST(vload_private),
-    ADD_TEST( vstore_global ),
+    ADD_TEST(vstore_global),
-    ADD_TEST( vstore_local ),
+    ADD_TEST(vstore_local),
-    ADD_TEST( vstore_private ),
+    ADD_TEST(vstore_private),
 
-    ADD_TEST( createkernelsinprogram ),
+    ADD_TEST(createkernelsinprogram),
-    ADD_TEST( imagedim_pow2 ),
+    ADD_TEST(imagedim_pow2),
-    ADD_TEST( imagedim_non_pow2 ),
+    ADD_TEST(imagedim_non_pow2),
-    ADD_TEST( image_param ),
+    ADD_TEST(image_param),
-    ADD_TEST( image_multipass_integer_coord ),
+    ADD_TEST(image_multipass_integer_coord),
-    ADD_TEST( image_multipass_float_coord ),
+    ADD_TEST(image_multipass_float_coord),
-    ADD_TEST( explicit_s2v_char ),
+    ADD_TEST(explicit_s2v_char),
-    ADD_TEST( explicit_s2v_uchar ),
+    ADD_TEST(explicit_s2v_uchar),
-    ADD_TEST( explicit_s2v_short ),
+    ADD_TEST(explicit_s2v_short),
-    ADD_TEST( explicit_s2v_ushort ),
+    ADD_TEST(explicit_s2v_ushort),
-    ADD_TEST( explicit_s2v_int ),
+    ADD_TEST(explicit_s2v_int),
-    ADD_TEST( explicit_s2v_uint ),
+    ADD_TEST(explicit_s2v_uint),
-    ADD_TEST( explicit_s2v_long ),
+    ADD_TEST(explicit_s2v_long),
-    ADD_TEST( explicit_s2v_ulong ),
+    ADD_TEST(explicit_s2v_ulong),
-    ADD_TEST( explicit_s2v_float ),
+    ADD_TEST(explicit_s2v_float),
-    ADD_TEST( explicit_s2v_double ),
+    ADD_TEST(explicit_s2v_double),
 
-    ADD_TEST( enqueue_map_buffer ),
+    ADD_TEST(enqueue_map_buffer),
-    ADD_TEST( enqueue_map_image ),
+    ADD_TEST(enqueue_map_image),
 
-    ADD_TEST( work_item_functions ),
+    ADD_TEST(work_item_functions),
 
-    ADD_TEST( astype ),
+    ADD_TEST(astype),
 
-    ADD_TEST( async_copy_global_to_local ),
+    ADD_TEST(async_copy_global_to_local),
-    ADD_TEST( async_copy_local_to_global ),
+    ADD_TEST(async_copy_local_to_global),
-    ADD_TEST( async_strided_copy_global_to_local ),
+    ADD_TEST(async_strided_copy_global_to_local),
-    ADD_TEST( async_strided_copy_local_to_global ),
+    ADD_TEST(async_strided_copy_local_to_global),
-    ADD_TEST( prefetch ),
+    ADD_TEST(prefetch),
 
-    ADD_TEST( kernel_call_kernel_function ),
+    ADD_TEST(kernel_call_kernel_function),
-    ADD_TEST( host_numeric_constants ),
+    ADD_TEST(host_numeric_constants),
-    ADD_TEST( kernel_numeric_constants ),
+    ADD_TEST(kernel_numeric_constants),
-    ADD_TEST( kernel_limit_constants ),
+    ADD_TEST(kernel_limit_constants),
-    ADD_TEST( kernel_preprocessor_macros ),
+    ADD_TEST(kernel_preprocessor_macros),
 
-    ADD_TEST( parameter_types ),
+    ADD_TEST(parameter_types),
-    ADD_TEST( vector_creation ),
+    ADD_TEST(vector_creation),
-    ADD_TEST( vec_type_hint ),
+    ADD_TEST(vector_swizzle),
-    ADD_TEST( kernel_memory_alignment_local ),
+    ADD_TEST(vec_type_hint),
-    ADD_TEST( kernel_memory_alignment_global ),
+    ADD_TEST(kernel_memory_alignment_local),
-    ADD_TEST( kernel_memory_alignment_constant ),
+    ADD_TEST(kernel_memory_alignment_global),
-    ADD_TEST( kernel_memory_alignment_private ),
+    ADD_TEST(kernel_memory_alignment_constant),
+    ADD_TEST(kernel_memory_alignment_private),
 
-    ADD_TEST_VERSION( progvar_prog_scope_misc, Version(2, 0) ),
+    ADD_TEST_VERSION(progvar_prog_scope_misc, Version(2, 0)),
-    ADD_TEST_VERSION( progvar_prog_scope_uninit, Version(2, 0) ),
+    ADD_TEST_VERSION(progvar_prog_scope_uninit, Version(2, 0)),
-    ADD_TEST_VERSION( progvar_prog_scope_init, Version(2, 0) ),
+    ADD_TEST_VERSION(progvar_prog_scope_init, Version(2, 0)),
-    ADD_TEST_VERSION( progvar_func_scope, Version(2, 0) ),
+    ADD_TEST_VERSION(progvar_func_scope, Version(2, 0)),
 
-    ADD_TEST( global_work_offsets ),
+    ADD_TEST(global_work_offsets),
-    ADD_TEST( get_global_offset ),
+    ADD_TEST(get_global_offset),
 
-    ADD_TEST_VERSION( global_linear_id, Version(2, 0) ),
+    ADD_TEST_VERSION(global_linear_id, Version(2, 0)),
-    ADD_TEST_VERSION( local_linear_id, Version(2, 0) ),
+    ADD_TEST_VERSION(local_linear_id, Version(2, 0)),
-    ADD_TEST_VERSION( enqueued_local_size, Version(2, 0) ),
+    ADD_TEST_VERSION(enqueued_local_size, Version(2, 0)),
 
-    ADD_TEST( simple_read_image_pitch ),
+    ADD_TEST(simple_read_image_pitch),
-    ADD_TEST( simple_write_image_pitch ),
+    ADD_TEST(simple_write_image_pitch),
 
 #if defined( __APPLE__ )
-    ADD_TEST( queue_priority ),
+    ADD_TEST(queue_priority),
 #endif
 
-    ADD_TEST_VERSION( get_linear_ids, Version(2, 0) ),
+    ADD_TEST_VERSION(get_linear_ids, Version(2, 0)),
-    ADD_TEST_VERSION( rw_image_access_qualifier, Version(2, 0) ),
+    ADD_TEST_VERSION(rw_image_access_qualifier, Version(2, 0)),
 };
 
 const int test_num = ARRAY_SIZE( test_list );

test_conformance/basic/procs.h

@@ -126,8 +126,12 @@ extern int    test_kernel_preprocessor_macros(cl_device_id deviceID, cl_context
 extern int      test_kernel_call_kernel_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
 
 extern int      test_parameter_types(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements);
-extern int      test_vector_creation(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
+extern int test_vector_creation(cl_device_id deviceID, cl_context context,
-extern int      test_vec_type_hint(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements);
+                                cl_command_queue queue, int num_elements);
+extern int test_vector_swizzle(cl_device_id deviceID, cl_context context,
+                               cl_command_queue queue, int num_elements);
+extern int test_vec_type_hint(cl_device_id deviceID, cl_context context,
+                              cl_command_queue queue, int num_elements);
 
 
 extern int test_kernel_memory_alignment_local(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems );

test_conformance/basic/test_arraycopy.cpp

@@ -51,7 +51,8 @@ test_arraycopy(cl_device_id device, cl_context context, cl_command_queue queue,
     output_ptr = (cl_uint*)malloc(sizeof(cl_uint) * num_elements);
 
     // results
-    results = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_uint) * num_elements, NULL, &err);
+    results = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                             sizeof(cl_uint) * num_elements, NULL, &err);
     test_error(err, "clCreateBuffer failed");
 
 /*****************************************************************************************************************************************/
@@ -64,7 +65,9 @@ test_arraycopy(cl_device_id device, cl_context context, cl_command_queue queue,
         input_ptr[i] = (cl_uint)(genrand_int32(d) & 0x7FFFFFFF);
 
     // client backing
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_USE_HOST_PTR), sizeof(cl_uint) * num_elements, input_ptr, &err);
+    streams[0] =
+        clCreateBuffer(context, CL_MEM_USE_HOST_PTR,
+                       sizeof(cl_uint) * num_elements, input_ptr, &err);
     test_error(err, "clCreateBuffer failed");
 
     delta_offset = num_elements * sizeof(cl_uint) / num_copies;
@@ -103,7 +106,8 @@ test_arraycopy(cl_device_id device, cl_context context, cl_command_queue queue,
         input_ptr[i] = (cl_uint)(genrand_int32(d) & 0x7FFFFFFF);
 
     // no backing
-    streams[2] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE) , sizeof(cl_uint) * num_elements, NULL, &err);
+    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_uint) * num_elements, NULL, &err);
     test_error(err, "clCreateBuffer failed");
 
     for (i=0; i<num_copies; i++)
@@ -146,17 +150,20 @@ test_arraycopy(cl_device_id device, cl_context context, cl_command_queue queue,
     free_mtdata(d); d= NULL;
 
     // client backing
-  streams[3] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_USE_HOST_PTR), sizeof(cl_uint) * num_elements, input_ptr, &err);
+    streams[3] =
-  test_error(err, "clCreateBuffer failed");
+        clCreateBuffer(context, CL_MEM_USE_HOST_PTR,
+                       sizeof(cl_uint) * num_elements, input_ptr, &err);
+    test_error(err, "clCreateBuffer failed");
 
-  err = create_single_kernel_helper(context, &program, &kernel, 1, &copy_kernel_code, "test_copy" );
+    err = create_single_kernel_helper(context, &program, &kernel, 1,
-  test_error(err, "create_single_kernel_helper failed");
+                                      &copy_kernel_code, "test_copy");
+    test_error(err, "create_single_kernel_helper failed");
 
-  err = clSetKernelArg(kernel, 0, sizeof streams[3], &streams[3]);
+    err = clSetKernelArg(kernel, 0, sizeof streams[3], &streams[3]);
-  err |= clSetKernelArg(kernel, 1, sizeof results, &results);
+    err |= clSetKernelArg(kernel, 1, sizeof results, &results);
-  test_error(err, "clSetKernelArg failed");
+    test_error(err, "clSetKernelArg failed");
 
-  size_t threads[3] = {num_elements, 0, 0};
+    size_t threads[3] = { num_elements, 0, 0 };
 
     err = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
   test_error(err, "clEnqueueNDRangeKernel failed");

test_conformance/basic/test_arrayimagecopy.cpp

@@ -38,7 +38,8 @@ int test_arrayimagecopy_single_format(cl_device_id device, cl_context context, c
 
   log_info("Testing %s %s\n", GetChannelOrderName(format->image_channel_order), GetChannelTypeName(format->image_channel_data_type));
 
-  image = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE), format, img_width, img_height, 0, NULL, &err);
+  image = create_image_2d(context, CL_MEM_READ_WRITE, format, img_width,
+                          img_height, 0, NULL, &err);
   test_error(err, "create_image_2d failed");
 
   err = clGetImageInfo(image, CL_IMAGE_ELEMENT_SIZE, sizeof(size_t), &elem_size, NULL);
@@ -46,7 +47,7 @@ int test_arrayimagecopy_single_format(cl_device_id device, cl_context context, c
 
   buffer_size = sizeof(cl_uchar) * elem_size * img_width * img_height;
 
-  buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  buffer_size, NULL, &err);
+  buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, buffer_size, NULL, &err);
   test_error(err, "clCreateBuffer failed");
 
   d = init_genrand( gRandomSeed );

test_conformance/basic/test_arrayimagecopy3d.cpp

@@ -39,7 +39,8 @@ int test_arrayimagecopy3d_single_format(cl_device_id device, cl_context context,
 
   log_info("Testing %s %s\n", GetChannelOrderName(format->image_channel_order), GetChannelTypeName(format->image_channel_data_type));
 
-  image = create_image_3d(context, (cl_mem_flags)(CL_MEM_READ_WRITE), format, img_width, img_height, img_depth, 0, 0, NULL, &err);
+  image = create_image_3d(context, CL_MEM_READ_ONLY, format, img_width,
+                          img_height, img_depth, 0, 0, NULL, &err);
   test_error(err, "create_image_3d failed");
 
   err = clGetImageInfo(image, CL_IMAGE_ELEMENT_SIZE, sizeof(size_t), &elem_size, NULL);
@@ -47,7 +48,7 @@ int test_arrayimagecopy3d_single_format(cl_device_id device, cl_context context,
 
   buffer_size = sizeof(cl_uchar) * elem_size * img_width * img_height * img_depth;
 
-  buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  buffer_size, NULL, &err);
+  buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, buffer_size, NULL, &err);
   test_error(err, "clCreateBuffer failed");
 
   d = init_genrand( gRandomSeed );
@@ -125,12 +126,15 @@ int test_arrayimagecopy3d(cl_device_id device, cl_context context, cl_command_qu
 
   PASSIVE_REQUIRE_3D_IMAGE_SUPPORT( device )
 
-  err = clGetSupportedImageFormats(context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE3D, 0, NULL, &num_formats);
+  err = clGetSupportedImageFormats(
+      context, CL_MEM_READ_ONLY, CL_MEM_OBJECT_IMAGE3D, 0, NULL, &num_formats);
   test_error(err, "clGetSupportedImageFormats failed");
 
   formats = (cl_image_format *)malloc(num_formats * sizeof(cl_image_format));
 
-  err = clGetSupportedImageFormats(context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE3D, num_formats, formats, NULL);
+  err = clGetSupportedImageFormats(context, CL_MEM_READ_ONLY,
+                                   CL_MEM_OBJECT_IMAGE3D, num_formats, formats,
+                                   NULL);
   test_error(err, "clGetSupportedImageFormats failed");
 
   for (i = 0; i < num_formats; i++) {

test_conformance/basic/test_arrayreadwrite.cpp

@@ -43,7 +43,8 @@ test_arrayreadwrite(cl_device_id device, cl_context context, cl_command_queue qu
     for (i=0; i<num_elements; i++)
         inptr[i] = (cl_uint)(genrand_int32(d) & 0x7FFFFFFF);
 
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_uint) * num_elements, NULL, &err);
+    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_uint) * num_elements, NULL, &err);
     test_error(err, "clCreateBuffer failed");
 
     for (i=0; i<num_tries; i++)

test_conformance/basic/test_barrier.cpp

@@ -108,11 +108,15 @@ test_barrier(cl_device_id device, cl_context context, cl_command_queue queue, in
     input_ptr = (int*)malloc(sizeof(int) * num_elements);
     output_ptr = (int*)malloc(sizeof(int));
 
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_int) * num_elements, NULL, &err);
+    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_int) * num_elements, NULL, &err);
     test_error(err, "clCreateBuffer failed.");
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_int), NULL, &err);
+    streams[1] =
+        clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int), NULL, &err);
     test_error(err, "clCreateBuffer failed.");
-    streams[2] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_int) * max_threadgroup_size, NULL, &err);
+    streams[2] =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       sizeof(cl_int) * max_threadgroup_size, NULL, &err);
     test_error(err, "clCreateBuffer failed.");
 
     d = init_genrand( gRandomSeed );

test_conformance/basic/test_constant.cpp

@@ -140,19 +140,22 @@ test_constant(cl_device_id device, cl_context context, cl_command_queue queue, i
     tmpI = (cl_int*)malloc(sizeof(cl_int) * constant_values);
     tmpF = (cl_float*)malloc(sizeof(cl_float) * constant_values);
     out  = (cl_float*)malloc(sizeof(cl_float) * constant_values);
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_float) * constant_values, NULL, NULL);
+    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_float) * constant_values, NULL, NULL);
     if (!streams[0])
     {
         log_error("clCreateBuffer failed\n");
         return -1;
     }
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_float) * constant_values, NULL, NULL);
+    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_float) * constant_values, NULL, NULL);
     if (!streams[1])
     {
         log_error("clCreateBuffer failed\n");
         return -1;
     }
-    streams[2] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_int) * constant_values, NULL, NULL);
+    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_int) * constant_values, NULL, NULL);
     if (!streams[2])
     {
         log_error("clCreateBuffer failed\n");

test_conformance/basic/test_enqueued_local_size.cpp

@@ -97,7 +97,8 @@ test_enqueued_local_size(cl_device_id device, cl_context context, cl_command_que
 
     output_ptr   = (int*)malloc(2 * sizeof(int));
 
-    streams = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), 2*sizeof(int), NULL, &err);
+    streams =
+        clCreateBuffer(context, CL_MEM_READ_WRITE, 2 * sizeof(int), NULL, &err);
     test_error( err, "clCreateBuffer failed.");
 
     std::string cl_std = "-cl-std=CL";

test_conformance/basic/test_explicit_s2v.cpp

@@ -146,9 +146,11 @@ int test_explicit_s2v_function(cl_device_id deviceID, cl_context context, cl_com
         return -1;
     }
 
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_COPY_HOST_PTR), paramSize * count, inputData, &error);
+    streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
+                                paramSize * count, inputData, &error);
     test_error( error, "clCreateBuffer failed");
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  destStride * count, NULL, &error);
+    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, destStride * count,
+                                NULL, &error);
     test_error( error, "clCreateBuffer failed");
 
     /* Set the arguments */

test_conformance/basic/test_float2int.cpp

@@ -69,13 +69,15 @@ test_float2int(cl_device_id device, cl_context context, cl_command_queue queue,
 
     input_ptr = (cl_float*)malloc(sizeof(cl_float) * num_elements);
     output_ptr = (cl_int*)malloc(sizeof(cl_int) * num_elements);
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_float) * num_elements, NULL, NULL);
+    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_float) * num_elements, NULL, NULL);
     if (!streams[0])
     {
         log_error("clCreateBuffer failed\n");
         return -1;
     }
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_int) * num_elements, NULL, NULL);
+    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_int) * num_elements, NULL, NULL);
     if (!streams[1])
     {
         log_error("clCreateBuffer failed\n");

test_conformance/basic/test_fpmath_float.cpp

@@ -161,13 +161,13 @@ test_fpmath_float(cl_device_id device, cl_context context, cl_command_queue queu
     input_ptr[2] = (cl_float*)malloc(length);
     output_ptr   = (cl_float*)malloc(length);
 
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length, NULL, &err);
+    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
     test_error( err, "clCreateBuffer failed.");
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length, NULL, &err);
+    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
     test_error( err, "clCreateBuffer failed.");
-    streams[2] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length, NULL, &err);
+    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
     test_error( err, "clCreateBuffer failed.");
-    streams[3] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length, NULL, &err);
+    streams[3] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
     test_error( err, "clCreateBuffer failed.");
 
     p = input_ptr[0];

test_conformance/basic/test_fpmath_float2.cpp

@@ -160,13 +160,13 @@ test_fpmath_float2(cl_device_id device, cl_context context, cl_command_queue que
     input_ptr[2] = (cl_float*)malloc(length);
     output_ptr   = (cl_float*)malloc(length);
 
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length, NULL, &err);
+    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
     test_error( err, "clCreateBuffer failed.");
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length, NULL, &err);
+    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
     test_error( err, "clCreateBuffer failed.");
-    streams[2] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length, NULL, &err);
+    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
     test_error( err, "clCreateBuffer failed.");
-    streams[3] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length, NULL, &err);
+    streams[3] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
     test_error( err, "clCreateBuffer failed.");
 
     p = input_ptr[0];

test_conformance/basic/test_fpmath_float4.cpp

@@ -160,13 +160,13 @@ test_fpmath_float4(cl_device_id device, cl_context context, cl_command_queue que
     input_ptr[2] = (cl_float*)malloc(length);
     output_ptr   = (cl_float*)malloc(length);
 
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length, NULL, &err);
+    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
     test_error( err, "clCreateBuffer failed.");
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length, NULL, &err);
+    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
     test_error( err, "clCreateBuffer failed.");
-    streams[2] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length, NULL, &err);
+    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
     test_error( err, "clCreateBuffer failed.");
-    streams[3] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length, NULL, &err);
+    streams[3] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
     test_error( err, "clCreateBuffer failed.");
 
     p = input_ptr[0];

test_conformance/basic/test_global_linear_id.cpp

@@ -70,14 +70,15 @@ test_global_linear_id(cl_device_id device, cl_context context, cl_command_queue
     cl_kernel kernel[2];
 
     int *output_ptr;
-      size_t threads[2];
+    size_t threads[2];
-      int err;
+    int err;
-      num_elements = (int)sqrt((float)num_elements);
+    num_elements = (int)sqrt((float)num_elements);
-      int length = num_elements * num_elements;
+    int length = num_elements * num_elements;
 
-      output_ptr   = (int*)malloc(sizeof(int) * length);
+    output_ptr = (int *)malloc(sizeof(int) * length);
 
-    streams = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length*sizeof(int), NULL, &err);
+    streams = clCreateBuffer(context, CL_MEM_READ_WRITE, length * sizeof(int),
+                             NULL, &err);
     test_error( err, "clCreateBuffer failed.");
 
     err = create_single_kernel_helper(context, &program[0], &kernel[0], 1,

test_conformance/basic/test_global_work_offsets.cpp

@@ -138,7 +138,9 @@ int test_global_work_offsets(cl_device_id deviceID, cl_context context, cl_comma
     memset( outputA, 0xff, sizeof( outputA ) );
     for( int i = 0; i < 3; i++ )
     {
-        streams[ i ] = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR), sizeof(outputA), outputA, &error );
+        streams[i] =
+            clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
+                           sizeof(outputA), outputA, &error);
         test_error( error, "Unable to create output array" );
     }
 
@@ -228,7 +230,9 @@ int test_get_global_offset(cl_device_id deviceID, cl_context context, cl_command
 
     // Create some output streams, and storage for a single control ID
     memset( outOffsets, 0xff, sizeof( outOffsets ) );
-    streams[0] = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR), sizeof( outOffsets ), outOffsets, &error );
+    streams[0] =
+        clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
+                       sizeof(outOffsets), outOffsets, &error);
     test_error( error, "Unable to create control ID buffer" );
 
     // Run a few different times

test_conformance/basic/test_hostptr.cpp

@@ -131,39 +131,53 @@ test_hostptr(cl_device_id device, cl_context context, cl_command_queue queue, in
         make_random_data(num_elements, input_ptr[1], d);
 
         // Create host-side input
-        streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_USE_HOST_PTR), sizeof(cl_float) * num_elements, input_ptr[0], &err);
+        streams[0] =
+            clCreateBuffer(context, CL_MEM_USE_HOST_PTR,
+                           sizeof(cl_float) * num_elements, input_ptr[0], &err);
         test_error(err, "clCreateBuffer 0 failed");
 
         // Create a copied input
-        streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_COPY_HOST_PTR), sizeof(cl_float) * num_elements, input_ptr[1], &err);
+        streams[1] =
+            clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
+                           sizeof(cl_float) * num_elements, input_ptr[1], &err);
         test_error(err, "clCreateBuffer 1 failed");
 
         // Create a host-side output
-        streams[2] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_USE_HOST_PTR), sizeof(cl_float) * num_elements, output_ptr, &err);
+        streams[2] =
+            clCreateBuffer(context, CL_MEM_USE_HOST_PTR,
+                           sizeof(cl_float) * num_elements, output_ptr, &err);
         test_error(err, "clCreateBuffer 2 failed");
 
         // Create a host-side input
         img_format.image_channel_order = CL_RGBA;
         img_format.image_channel_data_type = CL_UNORM_INT8;
-        streams[3] = create_image_2d(context, (cl_mem_flags)(CL_MEM_USE_HOST_PTR), &img_format, img_width, img_height, 0, rgba8_inptr, &err);
+        streams[3] =
+            create_image_2d(context, CL_MEM_USE_HOST_PTR, &img_format,
+                            img_width, img_height, 0, rgba8_inptr, &err);
         test_error(err, "create_image_2d 3 failed");
 
         // Create a copied input
         img_format.image_channel_order = CL_RGBA;
         img_format.image_channel_data_type = CL_UNORM_INT8;
-        streams[4] = create_image_2d(context, (cl_mem_flags)(CL_MEM_COPY_HOST_PTR), &img_format, img_width, img_height, 0, rgba8_inptr, &err);
+        streams[4] =
+            create_image_2d(context, CL_MEM_COPY_HOST_PTR, &img_format,
+                            img_width, img_height, 0, rgba8_inptr, &err);
         test_error(err, "create_image_2d 4 failed");
 
         // Create a host-side output
         img_format.image_channel_order = CL_RGBA;
         img_format.image_channel_data_type = CL_UNORM_INT8;
-        streams[5] = create_image_2d(context, (cl_mem_flags)(CL_MEM_USE_HOST_PTR), &img_format, img_width, img_height, 0, rgba8_outptr, &err);
+        streams[5] =
+            create_image_2d(context, CL_MEM_USE_HOST_PTR, &img_format,
+                            img_width, img_height, 0, rgba8_outptr, &err);
         test_error(err, "create_image_2d 5 failed");
 
         // Create a copied output
         img_format.image_channel_data_type = CL_RGBA;
         img_format.image_channel_data_type = CL_UNORM_INT8;
-        streams[6] = create_image_2d(context, (cl_mem_flags)(CL_MEM_COPY_HOST_PTR), &img_format, img_width, img_height, 0, rgba8_outptr, &err);
+        streams[6] =
+            create_image_2d(context, CL_MEM_COPY_HOST_PTR, &img_format,
+                            img_width, img_height, 0, rgba8_outptr, &err);
         test_error(err, "create_image_2d 6 failed");
 
         err = create_single_kernel_helper(context, &program, &kernel,1, &hostptr_kernel_code, "test_hostptr" );

test_conformance/basic/test_image_multipass.cpp

@@ -172,7 +172,7 @@ test_image_multipass_integer_coord(cl_device_id device, cl_context context, cl_c
         cl_mem_flags        flags;
 
         initial_data = generate_initial_byte_image(img_width, img_height, 4, 0xF0);
-        flags = (cl_mem_flags)(CL_MEM_READ_WRITE);
+        flags = CL_MEM_READ_WRITE;
 
         accum_streams[0] = create_image_2d(context, flags, &img_format, img_width, img_height, 0, NULL, NULL);
         if (!accum_streams[0])
@@ -224,7 +224,7 @@ test_image_multipass_integer_coord(cl_device_id device, cl_context context, cl_c
         MTdata              d;
 
         input_streams = (cl_mem*)malloc(sizeof(cl_mem) * num_input_streams);
-        flags = (cl_mem_flags)(CL_MEM_READ_WRITE);
+        flags = CL_MEM_READ_WRITE;
 
         int i;
         d = init_genrand( gRandomSeed );
@@ -425,7 +425,7 @@ test_image_multipass_float_coord(cl_device_id device, cl_context context, cl_com
         cl_mem_flags        flags;
 
         initial_data = generate_initial_byte_image(img_width, img_height, 4, 0xF0);
-        flags = (cl_mem_flags)(CL_MEM_READ_WRITE);
+        flags = CL_MEM_READ_WRITE;
 
         accum_streams[0] = create_image_2d(context, flags, &img_format, img_width, img_height, 0, NULL, NULL);
         if (!accum_streams[0])
@@ -469,7 +469,7 @@ test_image_multipass_float_coord(cl_device_id device, cl_context context, cl_com
         MTdata              d;
 
         input_streams = (cl_mem*)malloc(sizeof(cl_mem) * num_input_streams);
-        flags = (cl_mem_flags)(CL_MEM_READ_WRITE);
+        flags = CL_MEM_READ_WRITE;
 
         int i;
         d = init_genrand( gRandomSeed );

test_conformance/basic/test_image_r8.cpp

@@ -88,7 +88,9 @@ test_image_r8(cl_device_id device, cl_context context, cl_command_queue queue, i
     img_format.image_channel_data_type = CL_UNSIGNED_INT8;
 
     // early out if this image type is not supported
-    if( ! is_image_format_supported( context, (cl_mem_flags)(CL_MEM_READ_ONLY), CL_MEM_OBJECT_IMAGE2D, &img_format ) ) {
+    if (!is_image_format_supported(context, CL_MEM_READ_ONLY,
+                                   CL_MEM_OBJECT_IMAGE2D, &img_format))
+    {
         log_info("WARNING: Image type not supported; skipping test.\n");
         return 0;
     }
@@ -98,14 +100,17 @@ test_image_r8(cl_device_id device, cl_context context, cl_command_queue queue, i
     free_mtdata(d); d = NULL;
 
     output_ptr = (cl_uchar*)malloc(sizeof(cl_uchar) * img_width * img_height);
-    streams[0] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_ONLY), &img_format, img_width, img_height, 0, NULL, NULL);
+    streams[0] = create_image_2d(context, CL_MEM_READ_ONLY, &img_format,
+                                 img_width, img_height, 0, NULL, NULL);
     if (!streams[0])
     {
         log_error("create_image_2d failed\n");
         return -1;
     }
 
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_uchar) * img_width*img_height, NULL, NULL);
+    streams[1] =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       sizeof(cl_uchar) * img_width * img_height, NULL, NULL);
     if (!streams[1])
     {
         log_error("clCreateBuffer failed\n");

test_conformance/basic/test_imagearraycopy.cpp

@@ -38,7 +38,8 @@ int test_imagearraycopy_single_format(cl_device_id device, cl_context context, c
 
   log_info("Testing %s %s\n", GetChannelOrderName(format->image_channel_order), GetChannelTypeName(format->image_channel_data_type));
 
-  image = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE), format, img_width, img_height, 0, NULL, &err);
+  image = create_image_2d(context, CL_MEM_READ_WRITE, format, img_width,
+                          img_height, 0, NULL, &err);
   test_error(err, "create_image_2d failed");
 
   err = clGetImageInfo(image, CL_IMAGE_ELEMENT_SIZE, sizeof(size_t), &elem_size, NULL);
@@ -46,7 +47,7 @@ int test_imagearraycopy_single_format(cl_device_id device, cl_context context, c
 
   buffer_size = sizeof(cl_uchar) * elem_size * img_width * img_height;
 
-  buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  buffer_size, NULL, &err);
+  buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, buffer_size, NULL, &err);
   test_error(err, "clCreateBuffer failed");
 
   d = init_genrand( gRandomSeed );

test_conformance/basic/test_imagearraycopy3d.cpp

@@ -38,7 +38,8 @@ int test_imagearraycopy3d_single_format(cl_device_id device, cl_context context,
 
   log_info("Testing %s %s\n", GetChannelOrderName(format->image_channel_order), GetChannelTypeName(format->image_channel_data_type));
 
-  image = create_image_3d(context, (cl_mem_flags)(CL_MEM_READ_WRITE), format, img_width, img_height, img_depth, 0, 0, NULL, &err);
+  image = create_image_3d(context, CL_MEM_READ_ONLY, format, img_width,
+                          img_height, img_depth, 0, 0, NULL, &err);
   test_error(err, "create_image_3d failed");
 
   err = clGetImageInfo(image, CL_IMAGE_ELEMENT_SIZE, sizeof(size_t), &elem_size, NULL);
@@ -46,7 +47,7 @@ int test_imagearraycopy3d_single_format(cl_device_id device, cl_context context,
 
   buffer_size = sizeof(cl_uchar) * elem_size * img_width * img_height * img_depth;
 
-  buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  buffer_size, NULL, &err);
+  buffer = clCreateBuffer(context, CL_MEM_READ_WRITE, buffer_size, NULL, &err);
   test_error(err, "clCreateBuffer failed");
 
   d = init_genrand( gRandomSeed );
@@ -121,12 +122,15 @@ int test_imagearraycopy3d(cl_device_id device, cl_context context, cl_command_qu
 
   PASSIVE_REQUIRE_3D_IMAGE_SUPPORT( device )
 
-  err = clGetSupportedImageFormats(context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE3D, 0, NULL, &num_formats);
+  err = clGetSupportedImageFormats(
+      context, CL_MEM_READ_ONLY, CL_MEM_OBJECT_IMAGE3D, 0, NULL, &num_formats);
   test_error(err, "clGetSupportedImageFormats failed");
 
   formats = (cl_image_format *)malloc(num_formats * sizeof(cl_image_format));
 
-  err = clGetSupportedImageFormats(context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE3D, num_formats, formats, NULL);
+  err = clGetSupportedImageFormats(context, CL_MEM_READ_ONLY,
+                                   CL_MEM_OBJECT_IMAGE3D, num_formats, formats,
+                                   NULL);
   test_error(err, "clGetSupportedImageFormats failed");
 
   for (i = 0; i < num_formats; i++) {

test_conformance/basic/test_imagecopy.cpp

@@ -132,23 +132,29 @@ test_imagecopy(cl_device_id device, cl_context context, cl_command_queue queue,
 
     img_format.image_channel_order = CL_RGBA;
     img_format.image_channel_data_type = CL_UNORM_INT8;
-    streams[0] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, &err);
+    streams[0] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, &err);
     test_error(err, "create_image_2d failed");
-    streams[1] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, &err);
+    streams[1] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, &err);
     test_error(err, "create_image_2d failed");
 
     img_format.image_channel_order = CL_RGBA;
     img_format.image_channel_data_type = CL_UNORM_INT16;
-    streams[2] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, &err);
+    streams[2] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, &err);
     test_error(err, "create_image_2d failed");
-    streams[3] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, &err);
+    streams[3] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, &err);
     test_error(err, "create_image_2d failed");
 
     img_format.image_channel_order = CL_RGBA;
     img_format.image_channel_data_type = CL_FLOAT;
-    streams[4] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, &err);
+    streams[4] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, &err);
     test_error(err, "create_image_2d failed");
-    streams[5] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, &err);
+    streams[5] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, &err);
     test_error(err, "create_image_2d failed");
 
     for (i=0; i<3; i++)

test_conformance/basic/test_imagedim.cpp

@@ -163,7 +163,9 @@ test_imagedim_pow2(cl_device_id device, cl_context context, cl_command_queue que
 
             img_format.image_channel_order = CL_RGBA;
             img_format.image_channel_data_type = CL_UNORM_INT8;
-            streams[0] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, NULL);
+            streams[0] =
+                create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                img_width, img_height, 0, NULL, NULL);
             if (!streams[0])
             {
                 log_error("create_image_2d failed.  width = %d, height = %d\n", img_width, img_height);
@@ -174,7 +176,9 @@ test_imagedim_pow2(cl_device_id device, cl_context context, cl_command_queue que
             }
             img_format.image_channel_order = CL_RGBA;
             img_format.image_channel_data_type = CL_UNORM_INT8;
-            streams[1] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, NULL);
+            streams[1] =
+                create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                img_width, img_height, 0, NULL, NULL);
             if (!streams[1])
             {
                 log_error("create_image_2d failed.  width = %d, height = %d\n", img_width, img_height);
@@ -404,7 +408,9 @@ test_imagedim_non_pow2(cl_device_id device, cl_context context, cl_command_queue
 
                 img_format.image_channel_order = CL_RGBA;
                 img_format.image_channel_data_type = CL_UNORM_INT8;
-                streams[0] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, effective_img_width, effective_img_height, 0, NULL, NULL);
+                streams[0] = create_image_2d(
+                    context, CL_MEM_READ_WRITE, &img_format,
+                    effective_img_width, effective_img_height, 0, NULL, NULL);
                 if (!streams[0])
                 {
                     log_error("create_image_2d failed.  width = %d, height = %d\n", effective_img_width, effective_img_height);
@@ -415,7 +421,9 @@ test_imagedim_non_pow2(cl_device_id device, cl_context context, cl_command_queue
                 }
                 img_format.image_channel_order = CL_RGBA;
                 img_format.image_channel_data_type = CL_UNORM_INT8;
-                streams[1] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, effective_img_width, effective_img_height, 0, NULL, NULL);
+                streams[1] = create_image_2d(
+                    context, CL_MEM_READ_WRITE, &img_format,
+                    effective_img_width, effective_img_height, 0, NULL, NULL);
                 if (!streams[1])
                 {
                     log_error("create_image_2d failed.  width = %d, height = %d\n", effective_img_width, effective_img_height);

test_conformance/basic/test_imagenpot.cpp

@@ -110,8 +110,8 @@ test_imagenpot(cl_device_id device_id, cl_context context, cl_command_queue queu
 
         img_format.image_channel_order = CL_RGBA;
         img_format.image_channel_data_type = CL_UNORM_INT8;
-        streams[0] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format,
+        streams[0] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
-                                 img_width, img_height, 0, NULL, NULL);
+                                     img_width, img_height, 0, NULL, NULL);
         if (!streams[0])
         {
             log_error("create_image_2d failed\n");
@@ -120,8 +120,8 @@ test_imagenpot(cl_device_id device_id, cl_context context, cl_command_queue queu
         }
         img_format.image_channel_order = CL_RGBA;
         img_format.image_channel_data_type = CL_UNORM_INT8;
-    streams[1] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format,
+        streams[1] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
-                                 img_width, img_height, 0, NULL, NULL);
+                                     img_width, img_height, 0, NULL, NULL);
         if (!streams[1])
         {
             log_error("create_image_2d failed\n");

test_conformance/basic/test_imagerandomcopy.cpp

@@ -146,23 +146,29 @@ test_imagerandomcopy(cl_device_id device, cl_context context, cl_command_queue q
 
     img_format.image_channel_order = CL_RGBA;
     img_format.image_channel_data_type = CL_UNORM_INT8;
-    streams[0] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, &err);
+    streams[0] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, &err);
     test_error(err, "create_image_2d failed");
-    streams[1] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, &err);
+    streams[1] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, &err);
     test_error(err, "create_image_2d failed");
 
     img_format.image_channel_order = CL_RGBA;
     img_format.image_channel_data_type = CL_UNORM_INT16;
-    streams[2] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, &err);
+    streams[2] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, &err);
     test_error(err, "create_image_2d failed");
-    streams[3] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, &err);
+    streams[3] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, &err);
     test_error(err, "create_image_2d failed");
 
     img_format.image_channel_order = CL_RGBA;
     img_format.image_channel_data_type = CL_FLOAT;
-    streams[4] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, &err);
+    streams[4] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, &err);
     test_error(err, "create_image_2d failed");
-    streams[5] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, &err);
+    streams[5] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, &err);
     test_error(err, "create_image_2d failed");
 
     for (i=0; i<3; i++)

test_conformance/basic/test_imagereadwrite.cpp

@@ -215,17 +215,20 @@ test_imagereadwrite(cl_device_id device, cl_context context, cl_command_queue qu
 
     img_format.image_channel_order = CL_RGBA;
     img_format.image_channel_data_type = CL_UNORM_INT8;
-    streams[0] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, &err);
+    streams[0] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, &err);
     test_error(err, "create_image_2d failed");
 
     img_format.image_channel_order = CL_RGBA;
     img_format.image_channel_data_type = CL_UNORM_INT16;
-    streams[1] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, &err);
+    streams[1] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, &err);
     test_error(err, "create_image_2d failed");
 
     img_format.image_channel_order = CL_RGBA;
     img_format.image_channel_data_type = CL_FLOAT;
-    streams[2] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, &err);
+    streams[2] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, &err);
     test_error(err, "create_image_2d failed");
 
     for (i=0; i<3; i++)

test_conformance/basic/test_int2float.cpp

@@ -68,13 +68,15 @@ test_int2float(cl_device_id device, cl_context context, cl_command_queue queue,
 
     input_ptr = (cl_int*)malloc(sizeof(cl_int) * num_elements);
     output_ptr = (cl_float*)malloc(sizeof(cl_float) * num_elements);
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_int) * num_elements, NULL, NULL);
+    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_int) * num_elements, NULL, NULL);
     if (!streams[0])
     {
         log_error("clCreateBuffer failed\n");
         return -1;
     }
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_float) * num_elements, NULL, NULL);
+    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_float) * num_elements, NULL, NULL);
     if (!streams[1])
     {
         log_error("clCreateBuffer failed\n");

test_conformance/basic/test_local_linear_id.cpp

@@ -66,20 +66,21 @@ verify_local_linear_id(int *result, int n)
 int
 test_local_linear_id(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
 {
-      cl_mem streams;
+    cl_mem streams;
-      cl_program program[2];
+    cl_program program[2];
-      cl_kernel kernel[2];
+    cl_kernel kernel[2];
 
     int *output_ptr;
-      size_t threads[2];
+    size_t threads[2];
-      int err;
+    int err;
-      num_elements = (int)sqrt((float)num_elements);
+    num_elements = (int)sqrt((float)num_elements);
-      int length = num_elements * num_elements;
+    int length = num_elements * num_elements;
 
-      output_ptr   = (cl_int*)malloc(sizeof(int) * length);
+    output_ptr = (cl_int *)malloc(sizeof(int) * length);
 
-    streams = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length*sizeof(int), NULL, &err);
+    streams = clCreateBuffer(context, CL_MEM_READ_WRITE, length * sizeof(int),
-    test_error( err, "clCreateBuffer failed.");
+                             NULL, &err);
+    test_error(err, "clCreateBuffer failed.");
 
     err = create_single_kernel_helper(context, &program[0], &kernel[0], 1,
                                       &local_linear_id_1d_code,

test_conformance/basic/test_multireadimagemultifmt.cpp

@@ -136,7 +136,8 @@ test_mri_multiple(cl_device_id device, cl_context context, cl_command_queue queu
 
     img_format.image_channel_order = CL_RGBA;
     img_format.image_channel_data_type = CL_UNORM_INT8;
-    streams[0] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, NULL);
+    streams[0] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, NULL);
     if (!streams[0])
     {
         log_error("create_image_2d failed\n");
@@ -144,7 +145,8 @@ test_mri_multiple(cl_device_id device, cl_context context, cl_command_queue queu
     }
     img_format.image_channel_order = CL_RGBA;
     img_format.image_channel_data_type = CL_UNORM_INT16;
-    streams[1] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, NULL);
+    streams[1] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, NULL);
     if (!streams[1])
     {
         log_error("create_image_2d failed\n");
@@ -152,14 +154,17 @@ test_mri_multiple(cl_device_id device, cl_context context, cl_command_queue queu
     }
     img_format.image_channel_order = CL_RGBA;
     img_format.image_channel_data_type = CL_FLOAT;
-    streams[2] = create_image_2d(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  &img_format, img_width, img_height, 0, NULL, NULL);
+    streams[2] = create_image_2d(context, CL_MEM_READ_WRITE, &img_format,
+                                 img_width, img_height, 0, NULL, NULL);
     if (!streams[2])
     {
         log_error("create_image_2d failed\n");
         return -1;
     }
 
-    streams[3] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(float)*4 * img_width*img_height, NULL, NULL);
+    streams[3] =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       sizeof(float) * 4 * img_width * img_height, NULL, NULL);
     if (!streams[3])
     {
         log_error("clCreateBuffer failed\n");

test_conformance/basic/test_numeric_constants.cpp

@@ -242,11 +242,14 @@ int test_kernel_numeric_constants(cl_device_id deviceID, cl_context context, cl_
     }
 
     /* Create some I/O streams */
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(float_out), NULL, &error);
+    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float_out),
+                                NULL, &error);
     test_error( error, "Creating test array failed" );
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(int_out), NULL, &error);
+    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(int_out),
+                                NULL, &error);
     test_error( error, "Creating test array failed" );
-    streams[2] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(uint_out), NULL, &error);
+    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(uint_out),
+                                NULL, &error);
     test_error( error, "Creating test array failed" );
 
     error = clSetKernelArg(kernel, 1, sizeof( streams[1] ), &streams[1]);
@@ -348,9 +351,11 @@ int test_kernel_numeric_constants(cl_device_id deviceID, cl_context context, cl_
             return -1;
         }
 
-        streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(long_out), NULL, &error);
+        streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                    sizeof(long_out), NULL, &error);
         test_error( error, "Creating test array failed" );
-        streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(ulong_out), NULL, &error);
+        streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                    sizeof(ulong_out), NULL, &error);
         test_error( error, "Creating test array failed" );
 
         error = clSetKernelArg(kernel, 1, sizeof( streams[1] ), &streams[1]);
@@ -389,9 +394,11 @@ int test_kernel_numeric_constants(cl_device_id deviceID, cl_context context, cl_
             return -1;
         }
 
-        streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(double_out), NULL, &error);
+        streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                    sizeof(double_out), NULL, &error);
         test_error( error, "Creating test array failed" );
-        streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(long_out), NULL, &error);
+        streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                    sizeof(long_out), NULL, &error);
         test_error( error, "Creating test array failed" );
 
         error = clSetKernelArg(kernel, 1, sizeof( streams[1] ), &streams[1]);
@@ -533,9 +540,11 @@ int test_kernel_limit_constants(cl_device_id deviceID, cl_context context, cl_co
 
 
     /* Create some I/O streams */
-    intStream = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(intOut), NULL, &error );
+    intStream = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(intOut), NULL,
+                               &error);
     test_error( error, "Creating test array failed" );
-    floatStream = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(floatOut), NULL, &error );
+    floatStream = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(floatOut),
+                                 NULL, &error);
     test_error( error, "Creating test array failed" );
 
     // Stage 1: basic limits on MAXFLOAT
@@ -677,7 +686,8 @@ int test_kernel_limit_constants(cl_device_id deviceID, cl_context context, cl_co
                 return -1;
             }
 
-            doubleStream = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(doubleOut), NULL, &error );
+            doubleStream = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                          sizeof(doubleOut), NULL, &error);
             test_error( error, "Creating test array failed" );
 
             error = clSetKernelArg( kernel, 0, sizeof( intStream ), &intStream );

test_conformance/basic/test_preprocessors.cpp

@@ -125,11 +125,14 @@ int test_kernel_preprocessor_macros(cl_device_id deviceID, cl_context context, c
     }
 
     /* Create some I/O streams */
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(results), NULL, &error);
+    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(results),
+                                NULL, &error);
     test_error( error, "Creating test array failed" );
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(fileString), NULL, &error);
+    streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(fileString),
+                                NULL, &error);
     test_error( error, "Creating test array failed" );
-    streams[2] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(roundingString), NULL, &error);
+    streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(roundingString), NULL, &error);
     test_error( error, "Creating test array failed" );
 
     // Set up and run
@@ -213,33 +216,15 @@ int test_kernel_preprocessor_macros(cl_device_id deviceID, cl_context context, c
 
     // The OpenCL version reported by the macro reports the feature level supported by the compiler. Since
     // this doesn't directly match any property we can query, we just check to see if it's a sane value
-    char versionBuffer[ 128 ];
+    auto device_cl_version = get_device_cl_version(deviceID);
-    error = clGetDeviceInfo( deviceID, CL_DEVICE_VERSION, sizeof( versionBuffer ), versionBuffer, NULL );
+    int device_cl_version_int = device_cl_version.to_int() * 10;
-    test_error( error, "Unable to get device's version to validate against" );
+    if ((results[2] < 100) || (results[2] > device_cl_version_int))
-
-    // We need to parse to get the version number to compare against
-    char *p1, *p2, *p3;
-    for( p1 = versionBuffer; ( *p1 != 0 ) && !isdigit( *p1 ); p1++ )
-        ;
-    for( p2 = p1; ( *p2 != 0 ) && ( *p2 != '.' ); p2++ )
-        ;
-    for( p3 = p2; ( *p3 != 0 ) && ( *p3 != ' ' ); p3++ )
-        ;
-
-    if( p2 == p3 )
     {
-        log_error( "ERROR: Unable to verify OpenCL version string (platform string is incorrect format)\n" );
+        log_error("ERROR: Kernel preprocessor __OPENCL_VERSION__ does not make "
-        return -1;
+                  "sense w.r.t. device's version string! "
-    }
+                  "(preprocessor states %d, CL_DEVICE_VERSION is %d (%s))\n",
-    *p2 = 0;
+                  results[2], device_cl_version_int,
-    *p3 = 0;
+                  device_cl_version.to_string().c_str());
-    int major = atoi( p1 );
-    int minor = atoi( p2 + 1 );
-    int realVersion = ( major * 100 ) + ( minor * 10 );
-    if( ( results[ 2 ] < 100 ) || ( results[ 2 ] > realVersion ) )
-    {
-        log_error( "ERROR: Kernel preprocessor __OPENCL_VERSION__ does not make sense w.r.t. device's version string! "
-                  "(preprocessor states %d, real version is %d (%d.%d))\n", results[ 2 ], realVersion, major, minor );
         return -1;
     }
 
@@ -250,33 +235,29 @@ int test_kernel_preprocessor_macros(cl_device_id deviceID, cl_context context, c
         return -1;
     }
 
-    // The OpenCL C version reported by the macro reports the OpenCL C supported by the compiler for this OpenCL device.
+    // The OpenCL C version reported by the macro reports the OpenCL C version
-    char cVersionBuffer[ 128 ];
+    // specified to the compiler. We need to see whether it is supported.
-    error = clGetDeviceInfo( deviceID, CL_DEVICE_OPENCL_C_VERSION, sizeof( cVersionBuffer ), cVersionBuffer, NULL );
+    int cl_c_major_version = results[3] / 100;
-    test_error( error, "Unable to get device's OpenCL C version to validate against" );
+    int cl_c_minor_version = (results[3] / 10) % 10;
-
+    if ((results[3] < 100)
-    // We need to parse to get the version number to compare against
+        || (!device_supports_cl_c_version(
-    for( p1 = cVersionBuffer; ( *p1 != 0 ) && !isdigit( *p1 ); p1++ )
+            deviceID, Version{ cl_c_major_version, cl_c_minor_version })))
-        ;
-    for( p2 = p1; ( *p2 != 0 ) && ( *p2 != '.' ); p2++ )
-        ;
-    for( p3 = p2; ( *p3 != 0 ) && ( *p3 != ' ' ); p3++ )
-        ;
-
-    if( p2 == p3 )
     {
-        log_error( "ERROR: Unable to verify OpenCL C version string (platform string is incorrect format)\n" );
+        auto device_version = get_device_cl_c_version(deviceID);
-        return -1;
+        log_error(
-    }
+            "ERROR: Kernel preprocessor __OPENCL_C_VERSION__ does not make "
-    *p2 = 0;
+            "sense w.r.t. device's version string! "
-    *p3 = 0;
+            "(preprocessor states %d, CL_DEVICE_OPENCL_C_VERSION is %d (%s))\n",
-    major = atoi( p1 );
+            results[3], device_version.to_int() * 10,
-    minor = atoi( p2 + 1 );
+            device_version.to_string().c_str());
-    realVersion = ( major * 100 ) + ( minor * 10 );
+        log_error("This means that CL_DEVICE_OPENCL_C_VERSION < "
-    if( ( results[ 3 ] < 100 ) || ( results[ 3 ] > realVersion ) )
+                  "__OPENCL_C_VERSION__");
-    {
+        if (device_cl_version >= Version{ 3, 0 })
-        log_error( "ERROR: Kernel preprocessor __OPENCL_C_VERSION__ does not make sense w.r.t. device's version string! "
+        {
-                  "(preprocessor states %d, real version is %d (%d.%d))\n", results[ 2 ], realVersion, major, minor );
+            log_error(", and __OPENCL_C_VERSION__ does not appear in "
+                      "CL_DEVICE_OPENCL_C_ALL_VERSIONS");
+        }
+        log_error("\n");
         return -1;
     }
 

test_conformance/basic/test_queue_priority.cpp

@@ -235,18 +235,18 @@ int test_queue_priority(cl_device_id device, cl_context context, cl_command_queu
     oldMode = get_round();
   }
 
-    input_ptr[0] = (cl_float*)malloc(length);
+  input_ptr[0] = (cl_float *)malloc(length);
-    input_ptr[1] = (cl_float*)malloc(length);
+  input_ptr[1] = (cl_float *)malloc(length);
-    input_ptr[2] = (cl_float*)malloc(length);
+  input_ptr[2] = (cl_float *)malloc(length);
-    output_ptr   = (cl_float*)malloc(length);
+  output_ptr = (cl_float *)malloc(length);
 
-  streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length, NULL, &err);
+  streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
   test_error( err, "clCreateBuffer failed.");
-  streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length, NULL, &err);
+  streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
   test_error( err, "clCreateBuffer failed.");
-  streams[2] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length, NULL, &err);
+  streams[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
   test_error( err, "clCreateBuffer failed.");
-  streams[3] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), length, NULL, &err);
+  streams[3] = clCreateBuffer(context, CL_MEM_READ_WRITE, length, NULL, &err);
   test_error( err, "clCreateBuffer failed.");
 
   p = input_ptr[0];

test_conformance/basic/test_rw_image_access_qualifier.cpp

@@ -121,8 +121,7 @@ int test_rw_image_access_qualifier(cl_device_id device_id, cl_context context, c
     format.image_channel_data_type = CL_UNSIGNED_INT32;
 
     /* Create input image */
-    flags = (cl_mem_flags) (CL_MEM_READ_WRITE
+    flags = CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR;
-                            | CL_MEM_COPY_HOST_PTR);
     src_image = create_image_2d(context, flags, &format,
                                 size_x, size_y, 0,
                                 (void *)input, &err);

test_conformance/basic/test_vector_swizzle.cpp

@@ -0,0 +1,681 @@
+//
+// Copyright (c) 2020 The Khronos Group Inc.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+
+#include <algorithm>
+#include <numeric>
+#include <string>
+#include <vector>
+
+#include "procs.h"
+#include "harness/testHarness.h"
+
+template <int N> struct TestInfo
+{
+};
+
+template <> struct TestInfo<2>
+{
+    static const size_t vector_size = 2;
+
+    static constexpr const char* kernel_source_xyzw = R"CLC(
+__kernel void test_vector_swizzle_xyzw(TYPE value, __global TYPE* dst) {
+    int index = 0;
+
+    // lvalue swizzles
+    dst[index++].x = value.x;
+    dst[index++].y = value.x;
+    dst[index++].xy = value;
+    dst[index++].yx = value;
+
+    // rvalue swizzles
+    dst[index++] = value.x;
+    dst[index++] = value.y;
+    dst[index++] = value.xy;
+    dst[index++] = value.yx;
+}
+)CLC";
+
+    static constexpr const char* kernel_source_rgba = R"CLC(
+__kernel void test_vector_swizzle_rgba(TYPE value, __global TYPE* dst) {
+    int index = 0;
+
+    // lvalue swizzles
+    dst[index++].r = value.r;
+    dst[index++].g = value.r;
+    dst[index++].rg = value;
+    dst[index++].gr = value;
+
+    // rvalue swizzles
+    dst[index++] = value.r;
+    dst[index++] = value.g;
+    dst[index++] = value.rg;
+    dst[index++] = value.gr;
+}
+)CLC";
+
+    static constexpr const char* kernel_source_sN = R"CLC(
+__kernel void test_vector_swizzle_sN(TYPE value, __global TYPE* dst) {
+    int index = 0;
+
+    // lvalue swizzles
+    dst[index++].s0 = value.s0;
+    dst[index++].s1 = value.s0;
+    dst[index++].s01 = value;
+    dst[index++].s10 = value;
+
+    // rvalue swizzles
+    dst[index++] = value.s0;
+    dst[index++] = value.s1;
+    dst[index++] = value.s01;
+    dst[index++] = value.s10;
+}
+)CLC";
+};
+
+template <> struct TestInfo<3>
+{
+    static const size_t vector_size = 4; // sizeof(vec3) is four elements
+
+    static constexpr const char* kernel_source_xyzw = R"CLC(
+__kernel void test_vector_swizzle_xyzw(TYPE value, __global TYPE* dst) {
+    int index = 0;
+
+    // lvalue swizzles
+    dst[index++].x = value.x;
+    dst[index++].y = value.x;
+    dst[index++].z = value.x;
+    dst[index++].xyz = value;
+    dst[index++].zyx = value;
+
+    // rvalue swizzles
+    vstore3(value.x, 0, (__global BASETYPE*)(dst + index++));
+    vstore3(value.y, 0, (__global BASETYPE*)(dst + index++));
+    vstore3(value.z, 0, (__global BASETYPE*)(dst + index++));
+    vstore3(value.xyz, 0, (__global BASETYPE*)(dst + index++));
+    vstore3(value.zyx, 0, (__global BASETYPE*)(dst + index++));
+}
+)CLC";
+
+    static constexpr const char* kernel_source_rgba = R"CLC(
+__kernel void test_vector_swizzle_rgba(TYPE value, __global TYPE* dst) {
+    int index = 0;
+
+    // lvalue swizzles
+    dst[index++].r = value.r;
+    dst[index++].g = value.r;
+    dst[index++].b = value.r;
+    dst[index++].rgb = value;
+    dst[index++].bgr = value;
+
+    // rvalue swizzles
+    vstore3(value.r, 0, (__global BASETYPE*)(dst + index++));
+    vstore3(value.g, 0, (__global BASETYPE*)(dst + index++));
+    vstore3(value.b, 0, (__global BASETYPE*)(dst + index++));
+    vstore3(value.rgb, 0, (__global BASETYPE*)(dst + index++));
+    vstore3(value.bgr, 0, (__global BASETYPE*)(dst + index++));
+}
+)CLC";
+
+    static constexpr const char* kernel_source_sN = R"CLC(
+__kernel void test_vector_swizzle_sN(TYPE value, __global TYPE* dst) {
+    int index = 0;
+
+    // lvalue swizzles
+    dst[index++].s0 = value.s0;
+    dst[index++].s1 = value.s0;
+    dst[index++].s2 = value.s0;
+    dst[index++].s012 = value;
+    dst[index++].s210 = value;
+
+    // rvalue swizzles
+    vstore3(value.s0, 0, (__global BASETYPE*)(dst + index++));
+    vstore3(value.s1, 0, (__global BASETYPE*)(dst + index++));
+    vstore3(value.s2, 0, (__global BASETYPE*)(dst + index++));
+    vstore3(value.s012, 0, (__global BASETYPE*)(dst + index++));
+    vstore3(value.s210, 0, (__global BASETYPE*)(dst + index++));
+}
+)CLC";
+};
+
+template <> struct TestInfo<4>
+{
+    static const size_t vector_size = 4;
+
+    static constexpr const char* kernel_source_xyzw = R"CLC(
+__kernel void test_vector_swizzle_xyzw(TYPE value, __global TYPE* dst) {
+    int index = 0;
+
+    // lvalue swizzles
+    dst[index++].x = value.x;
+    dst[index++].y = value.x;
+    dst[index++].z = value.x;
+    dst[index++].w = value.x;
+    dst[index++].xyzw = value;
+    dst[index++].wzyx = value;
+
+    // rvalue swizzles
+    dst[index++] = value.x;
+    dst[index++] = value.y;
+    dst[index++] = value.z;
+    dst[index++] = value.w;
+    dst[index++] = value.xyzw;
+    dst[index++] = value.wzyx;
+}
+)CLC";
+
+    static constexpr const char* kernel_source_rgba = R"CLC(
+__kernel void test_vector_swizzle_rgba(TYPE value, __global TYPE* dst) {
+    int index = 0;
+
+    // lvalue swizzles
+    dst[index++].r = value.r;
+    dst[index++].g = value.r;
+    dst[index++].b = value.r;
+    dst[index++].a = value.r;
+    dst[index++].rgba = value;
+    dst[index++].abgr = value;
+
+    // rvalue swizzles
+    dst[index++] = value.r;
+    dst[index++] = value.g;
+    dst[index++] = value.b;
+    dst[index++] = value.a;
+    dst[index++] = value.rgba;
+    dst[index++] = value.abgr;
+}
+)CLC";
+
+    static constexpr const char* kernel_source_sN = R"CLC(
+__kernel void test_vector_swizzle_sN(TYPE value, __global TYPE* dst) {
+    int index = 0;
+
+    // lvalue swizzles
+    dst[index++].s0 = value.s0;
+    dst[index++].s1 = value.s0;
+    dst[index++].s2 = value.s0;
+    dst[index++].s3 = value.s0;
+    dst[index++].s0123 = value;
+    dst[index++].s3210 = value;
+
+    // rvalue swizzles
+    dst[index++] = value.s0;
+    dst[index++] = value.s1;
+    dst[index++] = value.s2;
+    dst[index++] = value.s3;
+    dst[index++] = value.s0123;
+    dst[index++] = value.s3210;
+}
+)CLC";
+};
+
+template <> struct TestInfo<8>
+{
+    static const size_t vector_size = 8;
+
+    static constexpr const char* kernel_source_xyzw = R"CLC(
+__kernel void test_vector_swizzle_xyzw(TYPE value, __global TYPE* dst) {
+    int index = 0;
+
+    // xwzw only for first four components!
+
+    // lvalue swizzles
+    dst[index++].x = value.x;
+    dst[index++].y = value.x;
+    dst[index++].z = value.x;
+    dst[index++].w = value.x;
+    dst[index++].s4 = value.s0;
+    dst[index++].s5 = value.s0;
+    dst[index++].s6 = value.s0;
+    dst[index++].s7 = value.s0;
+    dst[index].xyzw = value.s0123;
+    dst[index++].s4567 = value.s4567;
+    dst[index].s7654 = value.s0123;
+    dst[index++].wzyx = value.s4567;
+
+    // rvalue swizzles
+    dst[index++] = value.x;
+    dst[index++] = value.y;
+    dst[index++] = value.z;
+    dst[index++] = value.w;
+    dst[index++] = value.s4;
+    dst[index++] = value.s5;
+    dst[index++] = value.s6;
+    dst[index++] = value.s7;
+    dst[index++] = (TYPE)(value.xyzw, value.s4567);
+    dst[index++] = (TYPE)(value.s7654, value.wzyx);
+}
+)CLC";
+    static constexpr const char* kernel_source_rgba = R"CLC(
+__kernel void test_vector_swizzle_rgba(TYPE value, __global TYPE* dst) {
+    int index = 0;
+
+    // rgba only for first four components!
+
+    // lvalue swizzles
+    dst[index++].r = value.r;
+    dst[index++].g = value.r;
+    dst[index++].b = value.r;
+    dst[index++].a = value.r;
+    dst[index++].s4 = value.s0;
+    dst[index++].s5 = value.s0;
+    dst[index++].s6 = value.s0;
+    dst[index++].s7 = value.s0;
+    dst[index].rgba = value.s0123;
+    dst[index++].s4567 = value.s4567;
+    dst[index].s7654 = value.s0123;
+    dst[index++].abgr = value.s4567;
+
+    // rvalue swizzles
+    dst[index++] = value.r;
+    dst[index++] = value.g;
+    dst[index++] = value.b;
+    dst[index++] = value.a;
+    dst[index++] = value.s4;
+    dst[index++] = value.s5;
+    dst[index++] = value.s6;
+    dst[index++] = value.s7;
+    dst[index++] = (TYPE)(value.rgba, value.s4567);
+    dst[index++] = (TYPE)(value.s7654, value.abgr);
+}
+)CLC";
+    static constexpr const char* kernel_source_sN = R"CLC(
+__kernel void test_vector_swizzle_sN(TYPE value, __global TYPE* dst) {
+    int index = 0;
+
+    // lvalue swizzles
+    dst[index++].s0 = value.s0;
+    dst[index++].s1 = value.s0;
+    dst[index++].s2 = value.s0;
+    dst[index++].s3 = value.s0;
+    dst[index++].s4 = value.s0;
+    dst[index++].s5 = value.s0;
+    dst[index++].s6 = value.s0;
+    dst[index++].s7 = value.s0;
+    dst[index++].s01234567 = value;
+    dst[index++].s76543210 = value;
+
+    // rvalue swizzles
+    dst[index++] = value.s0;
+    dst[index++] = value.s1;
+    dst[index++] = value.s2;
+    dst[index++] = value.s3;
+    dst[index++] = value.s4;
+    dst[index++] = value.s5;
+    dst[index++] = value.s6;
+    dst[index++] = value.s7;
+    dst[index++] = value.s01234567;
+    dst[index++] = value.s76543210;
+}
+)CLC";
+};
+
+template <> struct TestInfo<16>
+{
+    static const size_t vector_size = 16;
+
+    static constexpr const char* kernel_source_xyzw = R"CLC(
+__kernel void test_vector_swizzle_xyzw(TYPE value, __global TYPE* dst) {
+    int index = 0;
+
+    // xwzw only for first four components!
+
+    // lvalue swizzles
+    dst[index++].x = value.x;
+    dst[index++].y = value.x;
+    dst[index++].z = value.x;
+    dst[index++].w = value.x;
+    dst[index++].s4 = value.s0;
+    dst[index++].s5 = value.s0;
+    dst[index++].s6 = value.s0;
+    dst[index++].s7 = value.s0;
+    dst[index++].s8 = value.s0;
+    dst[index++].s9 = value.s0;
+    dst[index++].sa = value.s0;
+    dst[index++].sb = value.s0;
+    dst[index++].sc = value.s0;
+    dst[index++].sd = value.s0;
+    dst[index++].se = value.s0;
+    dst[index++].sf = value.s0;
+    dst[index].xyzw = value.s0123;
+    dst[index].s4567 = value.s4567;
+    dst[index].s89ab = value.s89ab;
+    dst[index++].scdef = value.scdef;
+    dst[index].sfedc = value.s0123;
+    dst[index].sba98 = value.s4567;
+    dst[index].s7654 = value.s89ab;
+    dst[index++].wzyx = value.scdef;
+
+    // rvalue swizzles
+    dst[index++] = value.x;
+    dst[index++] = value.y;
+    dst[index++] = value.z;
+    dst[index++] = value.w;
+    dst[index++] = value.s4;
+    dst[index++] = value.s5;
+    dst[index++] = value.s6;
+    dst[index++] = value.s7;
+    dst[index++] = value.s8;
+    dst[index++] = value.s9;
+    dst[index++] = value.sa;
+    dst[index++] = value.sb;
+    dst[index++] = value.sc;
+    dst[index++] = value.sd;
+    dst[index++] = value.se;
+    dst[index++] = value.sf;
+    dst[index++] = (TYPE)(value.xyzw, value.s4567, value.s89abcdef);
+    dst[index++] = (TYPE)(value.sfedcba98, value.s7654, value.wzyx);
+}
+)CLC";
+    static constexpr const char* kernel_source_rgba = R"CLC(
+__kernel void test_vector_swizzle_rgba(TYPE value, __global TYPE* dst) {
+    int index = 0;
+
+    // rgba only for first four components!
+
+    // lvalue swizzles
+    dst[index++].r = value.r;
+    dst[index++].g = value.r;
+    dst[index++].b = value.r;
+    dst[index++].a = value.r;
+    dst[index++].s4 = value.s0;
+    dst[index++].s5 = value.s0;
+    dst[index++].s6 = value.s0;
+    dst[index++].s7 = value.s0;
+    dst[index++].s8 = value.s0;
+    dst[index++].s9 = value.s0;
+    dst[index++].sa = value.s0;
+    dst[index++].sb = value.s0;
+    dst[index++].sc = value.s0;
+    dst[index++].sd = value.s0;
+    dst[index++].se = value.s0;
+    dst[index++].sf = value.s0;
+    dst[index].rgba = value.s0123;
+    dst[index].s4567 = value.s4567;
+    dst[index].s89ab = value.s89ab;
+    dst[index++].scdef = value.scdef;
+    dst[index].sfedc = value.s0123;
+    dst[index].sba98 = value.s4567;
+    dst[index].s7654 = value.s89ab;
+    dst[index++].abgr = value.scdef;
+
+    // rvalue swizzles
+    dst[index++] = value.r;
+    dst[index++] = value.g;
+    dst[index++] = value.b;
+    dst[index++] = value.a;
+    dst[index++] = value.s4;
+    dst[index++] = value.s5;
+    dst[index++] = value.s6;
+    dst[index++] = value.s7;
+    dst[index++] = value.s8;
+    dst[index++] = value.s9;
+    dst[index++] = value.sa;
+    dst[index++] = value.sb;
+    dst[index++] = value.sc;
+    dst[index++] = value.sd;
+    dst[index++] = value.se;
+    dst[index++] = value.sf;
+    dst[index++] = (TYPE)(value.rgba, value.s4567, value.s89abcdef);
+    dst[index++] = (TYPE)(value.sfedcba98, value.s7654, value.abgr);
+}
+)CLC";
+    static constexpr const char* kernel_source_sN = R"CLC(
+__kernel void test_vector_swizzle_sN(TYPE value, __global TYPE* dst) {
+    int index = 0;
+
+    // lvalue swizzles
+    dst[index++].s0 = value.s0;
+    dst[index++].s1 = value.s0;
+    dst[index++].s2 = value.s0;
+    dst[index++].s3 = value.s0;
+    dst[index++].s4 = value.s0;
+    dst[index++].s5 = value.s0;
+    dst[index++].s6 = value.s0;
+    dst[index++].s7 = value.s0;
+    dst[index++].s8 = value.s0;
+    dst[index++].s9 = value.s0;
+    dst[index++].sa = value.s0;
+    dst[index++].sb = value.s0;
+    dst[index++].sc = value.s0;
+    dst[index++].sd = value.s0;
+    dst[index++].se = value.s0;
+    dst[index++].sf = value.s0;
+    dst[index++].s0123456789abcdef = value; // lower-case
+    dst[index++].sFEDCBA9876543210 = value; // upper-case
+
+    // rvalue swizzles
+    dst[index++] = value.s0;
+    dst[index++] = value.s1;
+    dst[index++] = value.s2;
+    dst[index++] = value.s3;
+    dst[index++] = value.s4;
+    dst[index++] = value.s5;
+    dst[index++] = value.s6;
+    dst[index++] = value.s7;
+    dst[index++] = value.s8;
+    dst[index++] = value.s9;
+    dst[index++] = value.sa;
+    dst[index++] = value.sb;
+    dst[index++] = value.sc;
+    dst[index++] = value.sd;
+    dst[index++] = value.se;
+    dst[index++] = value.sf;
+    dst[index++] = value.s0123456789abcdef; // lower-case
+    dst[index++] = value.sFEDCBA9876543210; // upper-case
+}
+)CLC";
+};
+
+template <typename T, size_t N, size_t S>
+static void makeReference(std::vector<T>& ref)
+{
+    // N single channel lvalue tests
+    // 2 multi-value lvalue tests
+    // N single channel rvalue tests
+    // 2 multi-value rvalue tests
+    const size_t refSize = (N + 2 + N + 2) * S;
+
+    ref.resize(refSize);
+    std::fill(ref.begin(), ref.end(), 99);
+
+    size_t dstIndex = 0;
+
+    // single channel lvalue
+    for (size_t i = 0; i < N; i++)
+    {
+        ref[dstIndex * S + i] = 0;
+        ++dstIndex;
+    }
+
+    // normal lvalue
+    for (size_t c = 0; c < N; c++)
+    {
+        ref[dstIndex * S + c] = c;
+    }
+    ++dstIndex;
+
+    // reverse lvalue
+    for (size_t c = 0; c < N; c++)
+    {
+        ref[dstIndex * S + c] = N - c - 1;
+    }
+    ++dstIndex;
+
+    // single channel rvalue
+    for (size_t i = 0; i < N; i++)
+    {
+        for (size_t c = 0; c < N; c++)
+        {
+            ref[dstIndex * S + c] = i;
+        }
+        ++dstIndex;
+    }
+
+    // normal rvalue
+    for (size_t c = 0; c < N; c++)
+    {
+        ref[dstIndex * S + c] = c;
+    }
+    ++dstIndex;
+
+    // reverse rvalue
+    for (size_t c = 0; c < N; c++)
+    {
+        ref[dstIndex * S + c] = N - c - 1;
+    }
+    ++dstIndex;
+
+    assert(dstIndex * S == refSize);
+}
+
+template <typename T>
+static int
+test_vectype_case(const std::vector<T>& value, const std::vector<T>& reference,
+                  cl_context context, cl_kernel kernel, cl_command_queue queue)
+{
+    cl_int error = CL_SUCCESS;
+
+    clMemWrapper mem;
+
+    std::vector<T> buffer(reference.size(), 99);
+    mem = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
+                         buffer.size() * sizeof(T), buffer.data(), &error);
+    test_error(error, "Unable to create test buffer");
+
+    error = clSetKernelArg(kernel, 0, value.size() * sizeof(T), value.data());
+    test_error(error, "Unable to set value kernel arg");
+
+    error = clSetKernelArg(kernel, 1, sizeof(mem), &mem);
+    test_error(error, "Unable to set destination buffer kernel arg");
+
+    size_t global_work_size[] = { 1 };
+    error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, global_work_size,
+                                   NULL, 0, NULL, NULL);
+    test_error(error, "Unable to enqueue test kernel");
+
+    error = clFinish(queue);
+    test_error(error, "clFinish failed after test kernel");
+
+    error =
+        clEnqueueReadBuffer(queue, mem, CL_TRUE, 0, buffer.size() * sizeof(T),
+                            buffer.data(), 0, NULL, NULL);
+    test_error(error, "Unable to read data after test kernel");
+
+    if (buffer != reference)
+    {
+        log_error("Result buffer did not match reference buffer!\n");
+        return TEST_FAIL;
+    }
+
+    return TEST_PASS;
+}
+
+template <typename T, size_t N>
+static int test_vectype(const char* type_name, cl_device_id device,
+                        cl_context context, cl_command_queue queue)
+{
+    log_info("    testing type %s%d\n", type_name, N);
+
+    cl_int error = CL_SUCCESS;
+    int result = TEST_PASS;
+
+    clProgramWrapper program;
+    clKernelWrapper kernel;
+
+    std::string buildOptions{ "-DTYPE=" };
+    buildOptions += type_name;
+    buildOptions += std::to_string(N);
+    buildOptions += " -DBASETYPE=";
+    buildOptions += type_name;
+
+    constexpr size_t S = TestInfo<N>::vector_size;
+
+    std::vector<T> value(S);
+    std::iota(value.begin(), value.end(), 0);
+
+    std::vector<T> reference;
+    makeReference<T, N, S>(reference);
+
+    // XYZW swizzles:
+
+    const char* xyzw_source = TestInfo<N>::kernel_source_xyzw;
+    error = create_single_kernel_helper(
+        context, &program, &kernel, 1, &xyzw_source, "test_vector_swizzle_xyzw",
+        buildOptions.c_str());
+    test_error(error, "Unable to create xyzw test kernel");
+
+    result |= test_vectype_case(value, reference, context, kernel, queue);
+
+    // sN swizzles:
+    const char* sN_source = TestInfo<N>::kernel_source_sN;
+    error = create_single_kernel_helper(context, &program, &kernel, 1,
+                                        &sN_source, "test_vector_swizzle_sN",
+                                        buildOptions.c_str());
+    test_error(error, "Unable to create sN test kernel");
+
+    result |= test_vectype_case(value, reference, context, kernel, queue);
+
+    // RGBA swizzles for OpenCL 3.0 and newer:
+    const Version device_version = get_device_cl_version(device);
+    if (device_version >= Version(3, 0))
+    {
+        const char* rgba_source = TestInfo<N>::kernel_source_rgba;
+        error = create_single_kernel_helper(
+            context, &program, &kernel, 1, &rgba_source,
+            "test_vector_swizzle_rgba", buildOptions.c_str());
+        test_error(error, "Unable to create rgba test kernel");
+
+        result |= test_vectype_case(value, reference, context, kernel, queue);
+    }
+
+    return result;
+}
+
+template <typename T>
+static int test_type(const char* type_name, cl_device_id device,
+                     cl_context context, cl_command_queue queue)
+{
+    return test_vectype<T, 2>(type_name, device, context, queue)
+        | test_vectype<T, 3>(type_name, device, context, queue)
+        | test_vectype<T, 4>(type_name, device, context, queue)
+        | test_vectype<T, 8>(type_name, device, context, queue)
+        | test_vectype<T, 16>(type_name, device, context, queue);
+}
+
+int test_vector_swizzle(cl_device_id device, cl_context context,
+                        cl_command_queue queue, int num_elements)
+{
+    int hasDouble = is_extension_available(device, "cl_khr_fp64");
+
+    int result = TEST_PASS;
+    result |= test_type<cl_char>("char", device, context, queue);
+    result |= test_type<cl_uchar>("uchar", device, context, queue);
+    result |= test_type<cl_short>("short", device, context, queue);
+    result |= test_type<cl_ushort>("ushort", device, context, queue);
+    result |= test_type<cl_int>("int", device, context, queue);
+    result |= test_type<cl_uint>("uint", device, context, queue);
+    if (gHasLong)
+    {
+        result |= test_type<cl_long>("long", device, context, queue);
+        result |= test_type<cl_ulong>("ulong", device, context, queue);
+    }
+    result |= test_type<cl_float>("float", device, context, queue);
+    if (hasDouble)
+    {
+        result |= test_type<cl_double>("double", device, context, queue);
+    }
+    return result;
+}

test_conformance/basic/test_wg_barrier.cpp

@@ -110,11 +110,15 @@ test_wg_barrier(cl_device_id device, cl_context context, cl_command_queue queue,
     input_ptr = (int*)malloc(sizeof(int) * num_elements);
     output_ptr = (int*)malloc(sizeof(int));
 
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_int) * num_elements, NULL, &err);
+    streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_int) * num_elements, NULL, &err);
     test_error(err, "clCreateBuffer failed.");
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_int), NULL, &err);
+    streams[1] =
+        clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int), NULL, &err);
     test_error(err, "clCreateBuffer failed.");
-    streams[2] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_int) * max_threadgroup_size, NULL, &err);
+    streams[2] =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       sizeof(cl_int) * max_threadgroup_size, NULL, &err);
     test_error(err, "clCreateBuffer failed.");
 
     d = init_genrand( gRandomSeed );

test_conformance/buffers/array_info.cpp

@@ -33,7 +33,8 @@ int test_array_info_size( cl_device_id deviceID, cl_context context, cl_command_
     size_t          retSize;
     size_t          elementSize = sizeof( cl_int );
 
-    memobj = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE),  elementSize * w*h*d, NULL, &err);
+    memobj = clCreateBuffer(context, CL_MEM_READ_WRITE, elementSize * w * h * d,
+                            NULL, &err);
     test_error(err, "clCreateBuffer failed.");
 
     err = clGetMemObjectInfo(memobj, CL_MEM_SIZE, sizeof( size_t ), (void *)&retSize, NULL);

test_conformance/buffers/test_buffer_mem.cpp

@@ -328,7 +328,8 @@ int test_mem_read_only_flags( cl_device_id deviceID, cl_context context, cl_comm
     for (i=0; i<num_elements; i++)
         inptr[i] = i;
 
-    buffers[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  sizeof(cl_int) * num_elements, NULL, &err);
+    buffers[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_int) * num_elements, NULL, &err);
     if ( err != CL_SUCCESS ){
         print_error(err, " clCreateBuffer failed to create MEM_ALLOC_GLOBAL_POOL array\n" );
         clReleaseMemObject( buffers[0]) ;

test_conformance/buffers/test_buffer_read.cpp

@@ -1233,7 +1233,8 @@ int test_buffer_read_struct(cl_device_id deviceID, cl_context context, cl_comman
         log_error( " unable to allocate %d bytes for output_ptr\n", (int)(objSize * num_elements) );
         return -1;
     }
-    buffers[0] = clCreateBuffer( context, (cl_mem_flags)(CL_MEM_READ_WRITE),  objSize * num_elements, NULL , &err);
+    buffers[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                objSize * num_elements, NULL, &err);
     if ( err != CL_SUCCESS ){
         print_error( err, " clCreateBuffer failed\n" );
         align_free( output_ptr );
@@ -1334,7 +1335,8 @@ static int testRandomReadSize( cl_device_id deviceID, cl_context context, cl_com
             }
             return -1;
         }
-        buffers[i] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE),  ptrSizes[i] * num_elements, NULL, &err);
+        buffers[i] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                    ptrSizes[i] * num_elements, NULL, &err);
         if ( err != CL_SUCCESS ){
             print_error(err, " clCreateBuffer failed\n" );
             for ( j = 0; j < i; j++ ){

test_conformance/buffers/test_buffer_write.cpp

@@ -554,8 +554,8 @@ static int verify_write_float( void *ptr1, void *ptr2, int n )
 static int verify_write_half( void *ptr1, void *ptr2, int n )
 {
     int     i;
-    cl_ushort   *inptr = (cl_ushort *)ptr1;
+    cl_half *inptr = (cl_half *)ptr1;
-    cl_ushort   *outptr = (cl_ushort *)ptr2;
+    cl_half *outptr = (cl_half *)ptr2;
 
     for ( i = 0; i < n; i++ ){
         if ( outptr[i] != inptr[i] )

test_conformance/c11_atomics/common.cpp

@@ -206,3 +206,77 @@ template<> cl_long AtomicTypeExtendedInfo<cl_long>::MaxValue() {return CL_LONG_M
 template<> cl_ulong AtomicTypeExtendedInfo<cl_ulong>::MaxValue() {return CL_ULONG_MAX;}
 template<> cl_float AtomicTypeExtendedInfo<cl_float>::MaxValue() {return CL_FLT_MAX;}
 template<> cl_double AtomicTypeExtendedInfo<cl_double>::MaxValue() {return CL_DBL_MAX;}
+
+cl_int getSupportedMemoryOrdersAndScopes(
+    cl_device_id device, std::vector<TExplicitMemoryOrderType> &memoryOrders,
+    std::vector<TExplicitMemoryScopeType> &memoryScopes)
+{
+    // The CL_DEVICE_ATOMIC_MEMORY_CAPABILITES is missing before 3.0, but since
+    // all orderings and scopes are required for 2.X devices and this test is
+    // skipped before 2.0 we can safely return all orderings and scopes if the
+    // device is 2.X. Query device for the supported orders.
+    if (get_device_cl_version(device) < Version{ 3, 0 })
+    {
+        memoryOrders.push_back(MEMORY_ORDER_EMPTY);
+        memoryOrders.push_back(MEMORY_ORDER_RELAXED);
+        memoryOrders.push_back(MEMORY_ORDER_ACQUIRE);
+        memoryOrders.push_back(MEMORY_ORDER_RELEASE);
+        memoryOrders.push_back(MEMORY_ORDER_ACQ_REL);
+        memoryOrders.push_back(MEMORY_ORDER_SEQ_CST);
+        memoryScopes.push_back(MEMORY_SCOPE_EMPTY);
+        memoryScopes.push_back(MEMORY_SCOPE_WORK_GROUP);
+        memoryScopes.push_back(MEMORY_SCOPE_DEVICE);
+        memoryScopes.push_back(MEMORY_SCOPE_ALL_SVM_DEVICES);
+        return CL_SUCCESS;
+    }
+
+    // For a 3.0 device we can query the supported orderings and scopes
+    // directly.
+    cl_device_atomic_capabilities atomic_capabilities{};
+    test_error(
+        clGetDeviceInfo(device, CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES,
+                        sizeof(atomic_capabilities), &atomic_capabilities,
+                        nullptr),
+        "clGetDeviceInfo failed for CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES\n");
+
+    // Provided we succeeded, we can start filling the vectors.
+    if (atomic_capabilities & CL_DEVICE_ATOMIC_ORDER_RELAXED)
+    {
+        memoryOrders.push_back(MEMORY_ORDER_RELAXED);
+    }
+
+    if (atomic_capabilities & CL_DEVICE_ATOMIC_ORDER_ACQ_REL)
+    {
+        memoryOrders.push_back(MEMORY_ORDER_ACQUIRE);
+        memoryOrders.push_back(MEMORY_ORDER_RELEASE);
+        memoryOrders.push_back(MEMORY_ORDER_ACQ_REL);
+    }
+
+    if (atomic_capabilities & CL_DEVICE_ATOMIC_ORDER_SEQ_CST)
+    {
+        // The functions not ending in explicit have the same semantics as the
+        // corresponding explicit function with memory_order_seq_cst for the
+        // memory_order argument.
+        memoryOrders.push_back(MEMORY_ORDER_EMPTY);
+        memoryOrders.push_back(MEMORY_ORDER_SEQ_CST);
+    }
+
+    if (atomic_capabilities & CL_DEVICE_ATOMIC_SCOPE_WORK_GROUP)
+    {
+        memoryScopes.push_back(MEMORY_SCOPE_WORK_GROUP);
+    }
+
+    if (atomic_capabilities & CL_DEVICE_ATOMIC_SCOPE_DEVICE)
+    {
+        // The functions that do not have memory_scope argument have the same
+        // semantics as the corresponding functions with the memory_scope
+        // argument set to memory_scope_device.
+        memoryScopes.push_back(MEMORY_SCOPE_EMPTY);
+        memoryScopes.push_back(MEMORY_SCOPE_DEVICE);
+    }
+    if (atomic_capabilities & CL_DEVICE_ATOMIC_SCOPE_ALL_DEVICES)
+    {
+        memoryScopes.push_back(MEMORY_SCOPE_ALL_SVM_DEVICES);
+    }
+    return CL_SUCCESS;
+}

test_conformance/c11_atomics/common.h

@@ -71,6 +71,10 @@ extern cl_device_atomic_capabilities gAtomicMemCap,
 extern const char *get_memory_order_type_name(TExplicitMemoryOrderType orderType);
 extern const char *get_memory_scope_type_name(TExplicitMemoryScopeType scopeType);
 
+extern cl_int getSupportedMemoryOrdersAndScopes(
+    cl_device_id device, std::vector<TExplicitMemoryOrderType> &memoryOrders,
+    std::vector<TExplicitMemoryScopeType> &memoryScopes);
+
 class AtomicTypeInfo
 {
 public:
@@ -487,16 +491,11 @@ public:
     std::vector<TExplicitMemoryScopeType> memoryScope;
     int error = 0;
 
-    memoryOrder.push_back(MEMORY_ORDER_EMPTY);
+    // For OpenCL-3.0 and later some orderings and scopes are optional, so here
-    memoryOrder.push_back(MEMORY_ORDER_RELAXED);
+    // we query for the supported ones.
-    memoryOrder.push_back(MEMORY_ORDER_ACQUIRE);
+    test_error_ret(
-    memoryOrder.push_back(MEMORY_ORDER_RELEASE);
+        getSupportedMemoryOrdersAndScopes(deviceID, memoryOrder, memoryScope),
-    memoryOrder.push_back(MEMORY_ORDER_ACQ_REL);
+        "getSupportedMemoryOrdersAndScopes failed\n", TEST_FAIL);
-    memoryOrder.push_back(MEMORY_ORDER_SEQ_CST);
-    memoryScope.push_back(MEMORY_SCOPE_EMPTY);
-    memoryScope.push_back(MEMORY_SCOPE_WORK_GROUP);
-    memoryScope.push_back(MEMORY_SCOPE_DEVICE);
-    memoryScope.push_back(MEMORY_SCOPE_ALL_SVM_DEVICES);
 
     for(unsigned oi = 0; oi < memoryOrder.size(); oi++)
     {
@@ -582,16 +581,11 @@ public:
     std::vector<TExplicitMemoryScopeType> memoryScope;
     int error = 0;
 
-    memoryOrder.push_back(MEMORY_ORDER_EMPTY);
+    // For OpenCL-3.0 and later some orderings and scopes are optional, so here
-    memoryOrder.push_back(MEMORY_ORDER_RELAXED);
+    // we query for the supported ones.
-    memoryOrder.push_back(MEMORY_ORDER_ACQUIRE);
+    test_error_ret(
-    memoryOrder.push_back(MEMORY_ORDER_RELEASE);
+        getSupportedMemoryOrdersAndScopes(deviceID, memoryOrder, memoryScope),
-    memoryOrder.push_back(MEMORY_ORDER_ACQ_REL);
+        "getSupportedMemoryOrdersAndScopes failed\n", TEST_FAIL);
-    memoryOrder.push_back(MEMORY_ORDER_SEQ_CST);
-    memoryScope.push_back(MEMORY_SCOPE_EMPTY);
-    memoryScope.push_back(MEMORY_SCOPE_WORK_GROUP);
-    memoryScope.push_back(MEMORY_SCOPE_DEVICE);
-    memoryScope.push_back(MEMORY_SCOPE_ALL_SVM_DEVICES);
 
     for(unsigned oi = 0; oi < memoryOrder.size(); oi++)
     {
@@ -800,23 +794,35 @@ std::string CBasicTest<HostAtomicType, HostDataType>::KernelCode(cl_uint maxNumD
     "\n";
   if(LocalMemory())
   {
-    code +=
+      // memory_order_relaxed is sufficient for these initialization operations
-      "  // initialize atomics not reachable from host (first thread is doing this, other threads are waiting on barrier)\n"
+      // as the barrier below will act as a fence, providing an order to the
-      "  if(get_local_id(0) == 0)\n"
+      // operations. memory_scope_work_group is sufficient as local memory is
-      "    for(uint dstItemIdx = 0; dstItemIdx < numDestItems; dstItemIdx++)\n"
+      // only visible within the work-group.
-      "    {\n";
+      code += R"(
-    if(aTypeName == "atomic_flag")
+              // initialize atomics not reachable from host (first thread
-    {
+              // is doing this, other threads are waiting on barrier)
-      code +=
+              if(get_local_id(0) == 0)
-        "      if(finalDest[dstItemIdx])\n"
+                for(uint dstItemIdx = 0; dstItemIdx < numDestItems; dstItemIdx++)
-        "        atomic_flag_test_and_set(destMemory+dstItemIdx);\n"
+                {)";
-        "      else\n"
+      if (aTypeName == "atomic_flag")
-        "        atomic_flag_clear(destMemory+dstItemIdx);\n";
+      {
-    }
+          code += R"(
+                  if(finalDest[dstItemIdx])
+                    atomic_flag_test_and_set_explicit(destMemory+dstItemIdx,
+                                                      memory_order_relaxed,
+                                                      memory_scope_work_group);
+                  else
+                    atomic_flag_clear_explicit(destMemory+dstItemIdx,
+                                               memory_order_relaxed,
+                                               memory_scope_work_group);)";
+      }
     else
     {
-      code +=
+        code += R"(
-        "      atomic_store(destMemory+dstItemIdx, finalDest[dstItemIdx]);\n";
+                atomic_store_explicit(destMemory+dstItemIdx,
+                                      finalDest[dstItemIdx],
+                                      memory_order_relaxed,
+                                      memory_scope_work_group);)";
     }
     code +=
       "    }\n"
@@ -873,20 +879,29 @@ std::string CBasicTest<HostAtomicType, HostDataType>::KernelCode(cl_uint maxNumD
         "  if(get_local_id(0) == 0) // first thread in workgroup\n";
     else
       // global atomics declared in program scope
-      code +=
+      code += R"(
-      "  if(atomic_fetch_add(&finishedThreads, 1) == get_global_size(0)-1)\n"
+                if(atomic_fetch_add_explicit(&finishedThreads, 1,
-      "    // last finished thread\n";
+                                           memory_order_relaxed,
+                                           memory_scope_work_group)
+                   == get_global_size(0)-1) // last finished thread
+                   )";
     code +=
         "    for(uint dstItemIdx = 0; dstItemIdx < numDestItems; dstItemIdx++)\n";
     if(aTypeName == "atomic_flag")
     {
-      code +=
+        code += R"(
-        "      finalDest[dstItemIdx] = atomic_flag_test_and_set(destMemory+dstItemIdx);\n";
+                finalDest[dstItemIdx] =
+                    atomic_flag_test_and_set_explicit(destMemory+dstItemIdx,
+                                                      memory_order_relaxed,
+                                                      memory_scope_work_group);)";
     }
     else
     {
-      code +=
+        code += R"(
-        "      finalDest[dstItemIdx] = atomic_load(destMemory+dstItemIdx);\n";
+                finalDest[dstItemIdx] =
+                    atomic_load_explicit(destMemory+dstItemIdx,
+                                         memory_order_relaxed,
+                                         memory_scope_work_group);)";
     }
   }
   code += "}\n"
@@ -948,51 +963,76 @@ int CBasicTest<HostAtomicType, HostDataType>::ExecuteSingleTest(cl_device_id dev
 
   if(deviceThreadCount > 0)
   {
-    cl_ulong usedLocalMemory;
+      // This loop iteratively reduces the workgroup size by 2 and then
-    cl_ulong totalLocalMemory;
+      // re-generates the kernel with the reduced
-    cl_uint maxWorkGroupSize;
+      // workgroup size until we find a size which is admissible for the kernel
+      // being run or reduce the wg size
+      // to the trivial case of 1 (which was separately verified to be accurate
+      // for the kernel being run)
 
-    // Set up the kernel code
+      while ((CurrentGroupSize() > 1))
-    programSource = PragmaHeader(deviceID)+ProgramHeader(numDestItems)+FunctionCode()+KernelCode(numDestItems);
+      {
-    programLine = programSource.c_str();
+          // Re-generate the kernel code with the current group size
-    if (create_single_kernel_helper_with_build_options(
+          if (kernel) clReleaseKernel(kernel);
-            context, &program, &kernel, 1, &programLine, "test_atomic_kernel",
+          if (program) clReleaseProgram(program);
-            gOldAPI ? "" : nullptr))
+          programSource = PragmaHeader(deviceID) + ProgramHeader(numDestItems)
-    {
+              + FunctionCode() + KernelCode(numDestItems);
-      return -1;
+          programLine = programSource.c_str();
-    }
+          if (create_single_kernel_helper_with_build_options(
-    if(gDebug)
+                  context, &program, &kernel, 1, &programLine,
-    {
+                  "test_atomic_kernel", gOldAPI ? "" : nullptr))
-      log_info("Program source:\n");
+          {
-      log_info("%s\n", programLine);
+              return -1;
-    }
+          }
-    // tune up work sizes based on kernel info
+          // Get work group size for the new kernel
-    error = clGetKernelWorkGroupInfo(kernel, deviceID, CL_KERNEL_WORK_GROUP_SIZE, sizeof(groupSize), &groupSize, NULL);
+          error = clGetKernelWorkGroupInfo(kernel, deviceID,
-    test_error(error, "Unable to obtain max work group size for device and kernel combo");
+                                           CL_KERNEL_WORK_GROUP_SIZE,
+                                           sizeof(groupSize), &groupSize, NULL);
+          test_error(error,
+                     "Unable to obtain max work group size for device and "
+                     "kernel combo");
 
-    if(LocalMemory())
+          if (LocalMemory())
-    {
+          {
-      error = clGetKernelWorkGroupInfo (kernel, deviceID, CL_KERNEL_LOCAL_MEM_SIZE, sizeof(usedLocalMemory), &usedLocalMemory, NULL);
+              cl_ulong usedLocalMemory;
-      test_error(error, "clGetKernelWorkGroupInfo failed");
+              cl_ulong totalLocalMemory;
+              cl_uint maxWorkGroupSize;
 
-      error = clGetDeviceInfo(deviceID, CL_DEVICE_LOCAL_MEM_SIZE, sizeof(totalLocalMemory), &totalLocalMemory, NULL);
+              error = clGetKernelWorkGroupInfo(
-      test_error(error, "clGetDeviceInfo failed");
+                  kernel, deviceID, CL_KERNEL_LOCAL_MEM_SIZE,
+                  sizeof(usedLocalMemory), &usedLocalMemory, NULL);
+              test_error(error, "clGetKernelWorkGroupInfo failed");
 
-      // We know that each work-group is going to use typeSize * deviceThreadCount bytes of local memory
+              error = clGetDeviceInfo(deviceID, CL_DEVICE_LOCAL_MEM_SIZE,
-      // so pick the maximum value for deviceThreadCount that uses all the local memory.
+                                      sizeof(totalLocalMemory),
-      maxWorkGroupSize = ((totalLocalMemory - usedLocalMemory) / typeSize);
+                                      &totalLocalMemory, NULL);
+              test_error(error, "clGetDeviceInfo failed");
 
-      if(maxWorkGroupSize < groupSize)
+              // We know that each work-group is going to use typeSize *
-        groupSize = maxWorkGroupSize;
+              // deviceThreadCount bytes of local memory
-    }
+              // so pick the maximum value for deviceThreadCount that uses all
+              // the local memory.
+              maxWorkGroupSize =
+                  ((totalLocalMemory - usedLocalMemory) / typeSize);
 
-    CurrentGroupSize((cl_uint)groupSize);
+              if (maxWorkGroupSize < groupSize) groupSize = maxWorkGroupSize;
+          }
+          if (CurrentGroupSize() <= groupSize)
+              break;
+          else
+              CurrentGroupSize(CurrentGroupSize() / 2);
+      }
     if(CurrentGroupSize() > deviceThreadCount)
       CurrentGroupSize(deviceThreadCount);
     if(CurrentGroupNum(deviceThreadCount) == 1 || gOldAPI)
       deviceThreadCount = CurrentGroupSize()*CurrentGroupNum(deviceThreadCount);
     threadCount = deviceThreadCount+hostThreadCount;
   }
+  if (gDebug)
+  {
+      log_info("Program source:\n");
+      log_info("%s\n", programLine);
+  }
   if(deviceThreadCount > 0)
     log_info("\t\t(thread count %u, group size %u)\n", deviceThreadCount, CurrentGroupSize());
   if(hostThreadCount > 0)
@@ -1037,11 +1077,13 @@ int CBasicTest<HostAtomicType, HostDataType>::ExecuteSingleTest(cl_device_id dev
       return -1;
     }
     memcpy(svmAtomicBuffer, &destItems[0], typeSize * numDestItems);
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_USE_HOST_PTR), typeSize * numDestItems, svmAtomicBuffer, NULL);
+    streams[0] = clCreateBuffer(context, CL_MEM_USE_HOST_PTR,
+                                typeSize * numDestItems, svmAtomicBuffer, NULL);
   }
   else
   {
-    streams[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_COPY_HOST_PTR), typeSize * numDestItems, &destItems[0], NULL);
+      streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
+                                  typeSize * numDestItems, &destItems[0], NULL);
   }
   if (!streams[0])
   {
@@ -1062,12 +1104,18 @@ int CBasicTest<HostAtomicType, HostDataType>::ExecuteSingleTest(cl_device_id dev
     }
     if(startRefValues.size())
       memcpy(svmDataBuffer, &startRefValues[0], typeSize*threadCount*NumNonAtomicVariablesPerThread());
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_USE_HOST_PTR), typeSize*threadCount*NumNonAtomicVariablesPerThread(), svmDataBuffer, NULL);
+    streams[1] = clCreateBuffer(context, CL_MEM_USE_HOST_PTR,
+                                typeSize * threadCount
+                                    * NumNonAtomicVariablesPerThread(),
+                                svmDataBuffer, NULL);
   }
   else
   {
-    streams[1] = clCreateBuffer(context, (cl_mem_flags)((startRefValues.size() ? CL_MEM_COPY_HOST_PTR : CL_MEM_READ_WRITE)),
+      streams[1] = clCreateBuffer(
-      typeSize * threadCount*NumNonAtomicVariablesPerThread(), startRefValues.size() ? &startRefValues[0] : 0, NULL);
+          context,
+          ((startRefValues.size() ? CL_MEM_COPY_HOST_PTR : CL_MEM_READ_WRITE)),
+          typeSize * threadCount * NumNonAtomicVariablesPerThread(),
+          startRefValues.size() ? &startRefValues[0] : 0, NULL);
   }
   if (!streams[1])
   {

test_conformance/c11_atomics/main.cpp

@@ -159,6 +159,32 @@ test_status InitCL(cl_device_id device) {
                 "Minimum atomic memory capabilities unsupported by device\n");
             return TEST_FAIL;
         }
+
+        // Disable program scope global variable testing in the case that it is
+        // not supported on an OpenCL-3.0 driver.
+        size_t max_global_variable_size{};
+        test_error_ret(clGetDeviceInfo(device,
+                                       CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE,
+                                       sizeof(max_global_variable_size),
+                                       &max_global_variable_size, nullptr),
+                       "Unable to get max global variable size\n", TEST_FAIL);
+        if (0 == max_global_variable_size)
+        {
+            gNoGlobalVariables = true;
+        }
+
+        // Disable generic address space testing in the case that it is not
+        // supported on an OpenCL-3.0 driver.
+        cl_bool generic_address_space_support{};
+        test_error_ret(
+            clGetDeviceInfo(device, CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT,
+                            sizeof(generic_address_space_support),
+                            &generic_address_space_support, nullptr),
+            "Unable to get generic address space support\n", TEST_FAIL);
+        if (CL_FALSE == generic_address_space_support)
+        {
+            gNoGenericAddressSpace = true;
+        }
     }
     else
     {

test_conformance/c11_atomics/test_atomics.cpp

@@ -206,6 +206,7 @@ public:
   using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrder;
   using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryScope;
   using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrderScopeStr;
+  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryScopeStr;
   using CBasicTest<HostAtomicType, HostDataType>::CheckCapabilities;
   CBasicTestLoad(TExplicitAtomicType dataType, bool useSVM) : CBasicTestMemOrderScope<HostAtomicType, HostDataType>(dataType, useSVM)
   {
@@ -228,11 +229,19 @@ public:
   }
   virtual std::string ProgramCore()
   {
-    std::string memoryOrderScope = MemoryOrderScopeStr();
+      // In the case this test is run with MEMORY_ORDER_ACQUIRE, the store
-    std::string postfix(memoryOrderScope.empty() ? "" : "_explicit");
+      // should be MEMORY_ORDER_RELEASE
-    return
+      std::string memoryOrderScopeLoad = MemoryOrderScopeStr();
-      "  atomic_store(&destMemory[tid], tid);\n"
+      std::string memoryOrderScopeStore =
-      "  oldValues[tid] = atomic_load"+postfix+"(&destMemory[tid]"+memoryOrderScope+");\n";
+          (MemoryOrder() == MEMORY_ORDER_ACQUIRE)
+          ? (", memory_order_release" + MemoryScopeStr())
+          : memoryOrderScopeLoad;
+      std::string postfix(memoryOrderScopeLoad.empty() ? "" : "_explicit");
+      return "  atomic_store" + postfix + "(&destMemory[tid], tid"
+          + memoryOrderScopeStore
+          + ");\n"
+            "  oldValues[tid] = atomic_load"
+          + postfix + "(&destMemory[tid]" + memoryOrderScopeLoad + ");\n";
   }
   virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
   {
@@ -1615,6 +1624,30 @@ public:
       orderStr = std::string(", ") + get_memory_order_type_name(MemoryOrderForClear());
     return orderStr + MemoryScopeStr();
   }
+
+  virtual int ExecuteSingleTest(cl_device_id deviceID, cl_context context,
+                                cl_command_queue queue)
+  {
+      // This test assumes support for the memory_scope_device scope in the case
+      // that LocalMemory() == false. Therefore we should skip this test in that
+      // configuration on a 3.0 driver since supporting the memory_scope_device
+      // scope is optionaly.
+      if (get_device_cl_version(deviceID) >= Version{ 3, 0 })
+      {
+          if (!LocalMemory()
+              && !(gAtomicFenceCap & CL_DEVICE_ATOMIC_SCOPE_DEVICE))
+          {
+              log_info(
+                  "Skipping atomic_flag test due to use of atomic_scope_device "
+                  "which is optionally not supported on this device\n");
+              return 0; // skip test - not applicable
+          }
+      }
+      return CBasicTestMemOrderScope<HostAtomicType,
+                                     HostDataType>::ExecuteSingleTest(deviceID,
+                                                                      context,
+                                                                      queue);
+  }
   virtual std::string ProgramCore()
   {
     std::string memoryOrderScope = MemoryOrderScopeStr();

test_conformance/clcpp/address_spaces/common.hpp

@@ -165,9 +165,8 @@ int run_address_spaces_test(cl_device_id device, cl_context context, cl_command_
     std::vector<TYPE> output = generate_output<TYPE>(work_size[0], 9999);
 
     // output buffer
-    buffers[0] = clCreateBuffer
+    buffers[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
-        (context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(TYPE) * output.size(), NULL, &err
+                                sizeof(TYPE) * output.size(), NULL, &err);
-    );
     RETURN_ON_CL_ERROR(err, "clCreateBuffer")
 
     // Execute test

test_conformance/clcpp/address_spaces/test_pointer_types.hpp

@@ -357,7 +357,8 @@ struct constant_pointer_test : public address_spaces_test<T>
         RETURN_ON_CL_ERROR(err, "clGetCommandQueueInfo");
 
         // Create constant buffer
-        auto const_buff = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_ONLY), sizeof(cl_uint), NULL, &err);
+        auto const_buff = clCreateBuffer(context, CL_MEM_READ_ONLY,
+                                         sizeof(cl_uint), NULL, &err);
         RETURN_ON_CL_ERROR(err, "clCreateBuffer");
 
         // Write m_test_value to const_buff

test_conformance/clcpp/api/test_ctors.hpp

@@ -128,7 +128,9 @@ int test_ctors_execution(cl_device_id device,
 
     // host vector, size == count, output[0...count-1] == 1
     std::vector<cl_uint> output(count, cl_uint(1));
-    output_buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * output.size(), NULL, &error);
+    output_buffer =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       sizeof(cl_uint) * output.size(), NULL, &error);
     RETURN_ON_CL_ERROR(error, "clCreateBuffer")
 
     error = clEnqueueWriteBuffer(queue, output_buffer, CL_TRUE, 0, sizeof(cl_uint) * output.size(), static_cast<void *>(output.data()), 0, NULL, NULL);
@@ -298,7 +300,9 @@ AUTO_TEST_CASE(test_global_scope_ctors_executed_once)
 
     // host vector, size == count, output[0...count-1] == 1
     std::vector<cl_uint> output(count, cl_uint(1));
-    output_buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * output.size(), NULL, &error);
+    output_buffer =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       sizeof(cl_uint) * output.size(), NULL, &error);
     RETURN_ON_CL_ERROR(error, "clCreateBuffer")
 
     for(size_t i = 0; i < 4; i++)
@@ -435,7 +439,9 @@ AUTO_TEST_CASE(test_global_scope_ctors_ndrange)
 
     // host vector, size == count, output[0...count-1] == 1
     std::vector<cl_uint> output(count, cl_uint(1));
-    output_buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * output.size(), NULL, &error);
+    output_buffer =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       sizeof(cl_uint) * output.size(), NULL, &error);
     RETURN_ON_CL_ERROR(error, "clCreateBuffer")
 
     error = clEnqueueWriteBuffer(

test_conformance/clcpp/api/test_dtors.hpp

@@ -114,7 +114,9 @@ AUTO_TEST_CASE(test_global_scope_dtor_is_executed)
     // host vector, size == count, output[0...count-1] == 0xbeefbeef (3203383023)
     // values in output __MUST BE__ greater than 0 for the test to work correctly
     std::vector<cl_uint> output(count, cl_uint(0xbeefbeef));
-    output_buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * output.size(), NULL, &error);
+    output_buffer =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       sizeof(cl_uint) * output.size(), NULL, &error);
     RETURN_ON_CL_ERROR(error, "clCreateBuffer")
 
     error = clEnqueueWriteBuffer(
@@ -297,7 +299,9 @@ AUTO_TEST_CASE(test_global_scope_dtors_executed_once)
     // values in output __MUST BE__ greater than 0 for the test to work correctly
     cl_uint init_value = cl_uint(0xbeefbeef);
     std::vector<cl_uint> output(count, init_value);
-    output_buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * output.size(), NULL, &error);
+    output_buffer =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       sizeof(cl_uint) * output.size(), NULL, &error);
     RETURN_ON_CL_ERROR(error, "clCreateBuffer")
 
     error = clEnqueueWriteBuffer(
@@ -497,7 +501,9 @@ AUTO_TEST_CASE(test_global_scope_dtor_ndrange)
     // host vector, size == count, output[0...count-1] == 0xbeefbeef (3203383023)
     // values in output __MUST BE__ greater than 0 for the test to work correctly
     std::vector<cl_uint> output(count, cl_uint(0xbeefbeef));
-    output_buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * output.size(), NULL, &error);
+    output_buffer =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       sizeof(cl_uint) * output.size(), NULL, &error);
     RETURN_ON_CL_ERROR(error, "clCreateBuffer")
 
     error = clEnqueueWriteBuffer(

test_conformance/clcpp/api/test_spec_consts.hpp

@@ -124,7 +124,9 @@ AUTO_TEST_CASE(test_spec_consts_defaults)
 
     // host vector, size == 1, output[0] == 1
     std::vector<cl_int> output(1, cl_int(1));
-    output_buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_int) * output.size(), NULL, &error);
+    output_buffer =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       sizeof(cl_int) * output.size(), NULL, &error);
     RETURN_ON_CL_ERROR(error, "clCreateBuffer")
 
     error = clEnqueueWriteBuffer(queue, output_buffer, CL_TRUE, 0, sizeof(cl_int) * output.size(), static_cast<void *>(output.data()), 0, NULL, NULL);
@@ -248,7 +250,9 @@ AUTO_TEST_CASE(test_spec_consts_many_constants)
 
     // host vector, size == 1, output[0] == 1
     std::vector<cl_int> output(1, cl_int(1));
-    output_buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_int) * output.size(), NULL, &error);
+    output_buffer =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       sizeof(cl_int) * output.size(), NULL, &error);
     RETURN_ON_CL_ERROR(error, "clCreateBuffer")
 
     error = clEnqueueWriteBuffer(queue, output_buffer, CL_TRUE, 0, sizeof(cl_int) * output.size(), static_cast<void *>(output.data()), 0, NULL, NULL);
@@ -443,7 +447,9 @@ AUTO_TEST_CASE(test_spec_consts_different_types)
 
     // host vector, size == 1, output[0] == 1
     std::vector<cl_int> output(1, cl_int(1));
-    output_buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_int) * output.size(), NULL, &error);
+    output_buffer =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       sizeof(cl_int) * output.size(), NULL, &error);
     RETURN_ON_CL_ERROR(error, "clCreateBuffer")
 
     error = clEnqueueWriteBuffer(queue, output_buffer, CL_TRUE, 0, sizeof(cl_int) * output.size(), static_cast<void *>(output.data()), 0, NULL, NULL);

test_conformance/clcpp/atomics/atomic_fetch.hpp

@@ -130,10 +130,12 @@ int test_atomic_fetch_func(cl_device_id device, cl_context context, cl_command_q
     std::vector<TYPE> input = generate_input<TYPE>(count, op.min1(), op.max1(), std::vector<TYPE>());
     std::vector<TYPE> output = generate_output<TYPE>((count - 1) / atomic_bucket_size + 1);
 
-    buffers[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(TYPE) * input.size(), NULL, &err);
+    buffers[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(TYPE) * input.size(), NULL, &err);
     RETURN_ON_CL_ERROR(err, "clCreateBuffer")
 
-    buffers[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(TYPE) * output.size(), NULL, &err);
+    buffers[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(TYPE) * output.size(), NULL, &err);
     RETURN_ON_CL_ERROR(err, "clCreateBuffer")
 
     err = clEnqueueWriteBuffer(

test_conformance/clcpp/math_funcs/logarithmic_funcs.hpp

@@ -49,7 +49,8 @@ int get_ilogb_nan_zero(cl_device_id device, cl_context context, cl_command_queue
 
     std::vector<cl_int> output = generate_output<cl_int>(2);
 
-    buffers[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_int) * output.size(), NULL, &err);
+    buffers[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_int) * output.size(), NULL, &err);
     RETURN_ON_CL_ERROR(err, "clCreateBuffer")
 
     err = clSetKernelArg(kernel, 0, sizeof(buffers[0]), &buffers[0]);

test_conformance/clcpp/program_scope_ctors_dtors/common.hpp

@@ -230,9 +230,8 @@ int run_ps_ctor_dtor_test(cl_device_id device, cl_context context, cl_command_qu
     std::vector<cl_uint> output = generate_output<cl_uint>(work_size[0], 9999);
 
     // device output buffer
-    buffers[0] = clCreateBuffer(
+    buffers[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
-        context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * output.size(), NULL, &err
+                                sizeof(cl_uint) * output.size(), NULL, &err);
-    );
     RETURN_ON_CL_ERROR(err, "clCreateBuffer")
 
     // Execute test

test_conformance/clcpp/spec_constants/common.hpp

@@ -213,9 +213,8 @@ int run_spec_constants_test(cl_device_id device, cl_context context, cl_command_
     std::vector<TYPE> output = generate_output<TYPE>(work_size[0], 9999);
 
     // device output buffer
-    buffers[0] = clCreateBuffer(
+    buffers[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
-        context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(TYPE) * output.size(), NULL, &err
+                                sizeof(TYPE) * output.size(), NULL, &err);
-    );
     RETURN_ON_CL_ERROR(err, "clCreateBuffer");
 
     // Execute test

test_conformance/clcpp/subgroups/test_sg_all.hpp

@@ -169,10 +169,12 @@ int sub_group_all(cl_device_id device, cl_context context, cl_command_queue queu
     std::vector<cl_uint> input = generate_input_sg_all(flat_work_size + 1, wg_size);
     std::vector<cl_uint> output = generate_output_sg_all(flat_work_size, wg_size);
 
-    buffers[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * input.size(), NULL, &err);
+    buffers[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_uint) * input.size(), NULL, &err);
     RETURN_ON_CL_ERROR(err, "clCreateBuffer");
 
-    buffers[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * output.size(), NULL, &err);
+    buffers[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_uint) * output.size(), NULL, &err);
     RETURN_ON_CL_ERROR(err, "clCreateBuffer");
 
     err = clEnqueueWriteBuffer(

test_conformance/clcpp/subgroups/test_sg_any.hpp

@@ -169,10 +169,12 @@ int sub_group_any(cl_device_id device, cl_context context, cl_command_queue queu
     std::vector<cl_uint> input = generate_input_sg_any(flat_work_size + 1, wg_size);
     std::vector<cl_uint> output = generate_output_sg_any(flat_work_size, wg_size);
 
-    buffers[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * input.size(), NULL, &err);
+    buffers[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_uint) * input.size(), NULL, &err);
     RETURN_ON_CL_ERROR(err, "clCreateBuffer");
 
-    buffers[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * output.size(), NULL, &err);
+    buffers[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_uint) * output.size(), NULL, &err);
     RETURN_ON_CL_ERROR(err, "clCreateBuffer");
 
     err = clEnqueueWriteBuffer(

test_conformance/clcpp/subgroups/test_sg_broadcast.hpp

@@ -156,10 +156,12 @@ int sub_group_broadcast(cl_device_id device, cl_context context, cl_command_queu
     std::vector<cl_uint> input = generate_input_sg_broadcast(flat_work_size, wg_size);
     std::vector<cl_uint> output = generate_output_sg_broadcast(flat_work_size, wg_size);
 
-    buffers[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * input.size(), NULL,&err);
+    buffers[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_uint) * input.size(), NULL, &err);
     RETURN_ON_CL_ERROR(err, "clCreateBuffer");
 
-    buffers[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * output.size(), NULL, &err);
+    buffers[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_uint) * output.size(), NULL, &err);
     RETURN_ON_CL_ERROR(err, "clCreateBuffer");
 
     err = clEnqueueWriteBuffer(

test_conformance/clcpp/subgroups/test_sg_reduce.hpp

@@ -223,10 +223,13 @@ int sub_group_reduce(cl_device_id device, cl_context context, cl_command_queue q
     std::vector<CL_INT_TYPE> input = generate_input<CL_INT_TYPE, op>(flat_work_size, wg_size);
     std::vector<CL_INT_TYPE> output = generate_output<CL_INT_TYPE, op>(flat_work_size, wg_size);
 
-    buffers[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(CL_INT_TYPE) * input.size(), NULL, &err);
+    buffers[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(CL_INT_TYPE) * input.size(), NULL, &err);
     RETURN_ON_CL_ERROR(err, "clCreateBuffer");
 
-    buffers[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(CL_INT_TYPE) * output.size(), NULL, &err);
+    buffers[1] =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       sizeof(CL_INT_TYPE) * output.size(), NULL, &err);
     RETURN_ON_CL_ERROR(err, "clCreateBuffer");
 
     err = clEnqueueWriteBuffer(

test_conformance/clcpp/subgroups/test_sg_scan_exclusive.hpp

@@ -210,10 +210,13 @@ int sub_group_scan_exclusive(cl_device_id device, cl_context context, cl_command
     std::vector<CL_INT_TYPE> input = generate_input<CL_INT_TYPE, op>(flat_work_size, wg_size);
     std::vector<CL_INT_TYPE> output = generate_output<CL_INT_TYPE, op>(flat_work_size, wg_size);
 
-    buffers[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(CL_INT_TYPE) * input.size(), NULL, &err);
+    buffers[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(CL_INT_TYPE) * input.size(), NULL, &err);
     RETURN_ON_CL_ERROR(err, "clCreateBuffer");
 
-    buffers[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(CL_INT_TYPE) * output.size(), NULL, &err);
+    buffers[1] =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       sizeof(CL_INT_TYPE) * output.size(), NULL, &err);
     RETURN_ON_CL_ERROR(err, "clCreateBuffer");
 
     err = clEnqueueWriteBuffer(

test_conformance/clcpp/subgroups/test_sg_scan_inclusive.hpp

@@ -210,10 +210,13 @@ int sub_group_scan_inclusive(cl_device_id device, cl_context context, cl_command
     std::vector<CL_INT_TYPE> input = generate_input<CL_INT_TYPE, op>(flat_work_size, wg_size);
     std::vector<CL_INT_TYPE> output = generate_output<CL_INT_TYPE, op>(flat_work_size, wg_size);
 
-    buffers[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(CL_INT_TYPE) * input.size(), NULL, &err);
+    buffers[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(CL_INT_TYPE) * input.size(), NULL, &err);
     RETURN_ON_CL_ERROR(err, "clCreateBuffer");
 
-    buffers[1] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(CL_INT_TYPE) * output.size(), NULL, &err);
+    buffers[1] =
+        clCreateBuffer(context, CL_MEM_READ_WRITE,
+                       sizeof(CL_INT_TYPE) * output.size(), NULL, &err);
     RETURN_ON_CL_ERROR(err, "clCreateBuffer");
 
     err = clEnqueueWriteBuffer(

test_conformance/clcpp/synchronization/named_barrier/common.hpp

@@ -133,7 +133,8 @@ int run_work_group_named_barrier_barrier_test(cl_device_id device, cl_context co
     std::vector<cl_uint> output = generate_output<cl_uint>(work_size[0], 9999);
 
     // device output buffer
-    buffers[0] = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * output.size(), NULL, &err);
+    buffers[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
+                                sizeof(cl_uint) * output.size(), NULL, &err);
     RETURN_ON_CL_ERROR(err, "clCreateBuffer")
 
     // Execute test kernels

test_conformance/clcpp/synchronization/named_barrier/test_named_barrier.hpp

@@ -276,7 +276,9 @@ struct global_fence_named_barrier_test : public work_group_named_barrier_test_ba
         RETURN_ON_CL_ERROR(err, "clGetCommandQueueInfo")
 
         // create temp buffer
-        auto temp_buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * work_size, NULL, &err);
+        auto temp_buffer =
+            clCreateBuffer(context, CL_MEM_READ_WRITE,
+                           sizeof(cl_uint) * work_size, NULL, &err);
         RETURN_ON_CL_ERROR(err, "clCreateBuffer")
 
         err = clSetKernelArg(kernel, 0, sizeof(output_buffer), &output_buffer);
@@ -420,10 +422,9 @@ struct global_local_fence_named_barrier_test : public work_group_named_barrier_t
         RETURN_ON_CL_ERROR(err, "clGetCommandQueueInfo")
 
         // create temp buffer
-        auto temp_buffer = clCreateBuffer(
+        auto temp_buffer =
-            context, (cl_mem_flags)(CL_MEM_READ_WRITE),
+            clCreateBuffer(context, CL_MEM_READ_WRITE,
-            sizeof(cl_uint) * work_size, NULL, &err
+                           sizeof(cl_uint) * work_size, NULL, &err);
-        );
         RETURN_ON_CL_ERROR(err, "clCreateBuffer")
 
         err = clSetKernelArg(kernel, 0, sizeof(output_buffer), &output_buffer);

test_conformance/clcpp/synchronization/named_barrier/test_spec_example.hpp

@@ -270,7 +270,9 @@ struct spec_example_work_group_named_barrier_test : public work_group_named_barr
         RETURN_ON_CL_ERROR(err, "clGetCommandQueueInfo")
 
         // create temp buffer
-        auto temp_buffer = clCreateBuffer(context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(cl_uint) * work_size, NULL, &err);
+        auto temp_buffer =
+            clCreateBuffer(context, CL_MEM_READ_WRITE,
+                           sizeof(cl_uint) * work_size, NULL, &err);
         RETURN_ON_CL_ERROR(err, "clCreateBuffer")
 
         err = clSetKernelArg(kernel, 0, sizeof(output_buffer), &output_buffer);

test_conformance/clcpp/utils_test/binary.hpp

@@ -243,19 +243,16 @@ int test_binary_func(cl_device_id device, cl_context context, cl_command_queue q
     std::vector<INPUT2> input2 = generate_input<INPUT2>(count, op.min2(), op.max2(), in2_spec_cases);
     std::vector<OUTPUT> output = generate_output<OUTPUT>(count);
 
-    buffers[0] = clCreateBuffer(
+    buffers[0] = clCreateBuffer(context, CL_MEM_READ_WRITE,
-        context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(INPUT1) * input1.size(), NULL, &err
+                                sizeof(INPUT1) * input1.size(), NULL, &err);
-    );
     RETURN_ON_CL_ERROR(err, "clCreateBuffer")
 
-    buffers[1] = clCreateBuffer(
+    buffers[1] = clCreateBuffer(context, CL_MEM_READ_WRITE,
-        context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(INPUT2) * input2.size(), NULL, &err
+                                sizeof(INPUT2) * input2.size(), NULL, &err);
-    );
     RETURN_ON_CL_ERROR(err, "clCreateBuffer")
 
-    buffers[2] = clCreateBuffer(
+    buffers[2] = clCreateBuffer(context, CL_MEM_READ_WRITE,
-        context, (cl_mem_flags)(CL_MEM_READ_WRITE), sizeof(OUTPUT) * output.size(), NULL, &err
+                                sizeof(OUTPUT) * output.size(), NULL, &err);
-    );
     RETURN_ON_CL_ERROR(err, "clCreateBuffer")
 
     err = clEnqueueWriteBuffer(