Removed usage of half types in CTS test KhronosGroup#1982

Author: kamil-goras-mobica

test_conformance/gl/common.h

@@ -51,7 +51,6 @@ static const format common_formats[] = {
   { GL_RGBA16UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_SHORT,               kUShort },
   { GL_RGBA32UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_INT,                 kUInt },
   { GL_RGBA32F_ARB,  GL_RGBA,             GL_FLOAT,                        kFloat },
-  { GL_RGBA16F_ARB,  GL_RGBA,             GL_HALF_FLOAT,                   kHalf }
 };
 
 #ifdef GL_VERSION_3_2

test_conformance/gl/test_images_write_common.cpp

@@ -36,14 +36,6 @@ static const char *kernelpattern_image_write_1D =
 "    write_image%s( dest, index, %s(value));\n"
 "}\n";
 
-static const char *kernelpattern_image_write_1D_half =
-"#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"
-"__kernel void sample_test( __global half4 *source, write_only image1d_t dest )\n"
-"{\n"
-"    uint index = get_global_id(0);\n"
-"    write_imagef( dest, index, vload_half4(index, (__global half *)source));\n"
-"}\n";
-
 static const char *kernelpattern_image_write_1D_buffer =
 "__kernel void sample_test( __global %s4 *source, write_only image1d_buffer_t dest )\n"
 "{\n"
@@ -52,14 +44,6 @@ static const char *kernelpattern_image_write_1D_buffer =
 "    write_image%s( dest, index, %s(value));\n"
 "}\n";
 
-static const char *kernelpattern_image_write_1D_buffer_half =
-"#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"
-"__kernel void sample_test( __global half4 *source, write_only image1d_buffer_t dest )\n"
-"{\n"
-"    uint index = get_global_id(0);\n"
-"    write_imagef( dest, index, vload_half4(index, (__global half *)source));\n"
-"}\n";
-
 static const char *kernelpattern_image_write_2D =
 "__kernel void sample_test( __global %s4 *source, write_only image2d_t dest )\n"
 "{\n"
@@ -70,16 +54,6 @@ static const char *kernelpattern_image_write_2D =
 "    write_image%s( dest, (int2)( tidX, tidY ), %s(value));\n"
 "}\n";
 
-static const char *kernelpattern_image_write_2D_half =
-"#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"
-"__kernel void sample_test( __global half4 *source, write_only image2d_t dest )\n"
-"{\n"
-"    int  tidX = get_global_id(0);\n"
-"    int  tidY = get_global_id(1);\n"
-"    uint index = tidY * get_image_width( dest ) + tidX;\n"
-"    write_imagef( dest, (int2)( tidX, tidY ), vload_half4(index, (__global half *)source));\n"
-"}\n";
-
 static const char *kernelpattern_image_write_1Darray =
 "__kernel void sample_test( __global %s4 *source, write_only image1d_array_t dest )\n"
 "{\n"
@@ -90,16 +64,6 @@ static const char *kernelpattern_image_write_1Darray =
 "    write_image%s( dest, (int2)( tidX, tidY ), %s(value));\n"
 "}\n";
 
-static const char *kernelpattern_image_write_1Darray_half =
-"#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"
-"__kernel void sample_test( __global half4 *source, write_only image1d_array_t dest )\n"
-"{\n"
-"    int  tidX = get_global_id(0);\n"
-"    int  tidY = get_global_id(1);\n"
-"    uint index = tidY * get_image_width( dest ) + tidX;\n"
-"    write_imagef( dest, (int2)( tidX, tidY ), vload_half4(index, (__global half *)source));\n"
-"}\n";
-
 static const char *kernelpattern_image_write_3D =
 "#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable\n"
 "__kernel void sample_test( __global %s4 *source, write_only image3d_t dest )\n"
@@ -114,20 +78,6 @@ static const char *kernelpattern_image_write_3D =
 "    write_image%s( dest, (int4)( tidX, tidY, tidZ, 0 ), %s(value));\n"
 "}\n";
 
-static const char *kernelpattern_image_write_3D_half =
-"#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"
-"#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable\n"
-"__kernel void sample_test( __global half4 *source, write_only image3d_t dest )\n"
-"{\n"
-"    int  tidX   = get_global_id(0);\n"
-"    int  tidY   = get_global_id(1);\n"
-"    int  tidZ   = get_global_id(2);\n"
-"    int  width  = get_image_width( dest );\n"
-"    int  height = get_image_height( dest );\n"
-"    int  index = tidZ * width * height + tidY * width + tidX;\n"
-"    write_imagef( dest, (int4)( tidX, tidY, tidZ, 0 ), vload_half4(index, (__global half *)source));\n"
-"}\n";
-
 static const char *kernelpattern_image_write_2Darray =
 "__kernel void sample_test( __global %s4 *source, write_only image2d_array_t dest )\n"
 "{\n"
@@ -141,19 +91,6 @@ static const char *kernelpattern_image_write_2Darray =
 "    write_image%s( dest, (int4)( tidX, tidY, tidZ, 0 ), %s(value));\n"
 "}\n";
 
-static const char *kernelpattern_image_write_2Darray_half =
-"#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n"
-"__kernel void sample_test( __global half4 *source, write_only image2d_array_t dest )\n"
-"{\n"
-"    int  tidX   = get_global_id(0);\n"
-"    int  tidY   = get_global_id(1);\n"
-"    int  tidZ   = get_global_id(2);\n"
-"    int  width  = get_image_width( dest );\n"
-"    int  height = get_image_height( dest );\n"
-"    int  index = tidZ * width * height + tidY * width + tidX;\n"
-"    write_imagef( dest, (int4)( tidX, tidY, tidZ, 0 ), vload_half4(index, (__global half *)source));\n"
-"}\n";
-
 #ifdef GL_VERSION_3_2
 
 static const char * kernelpattern_image_write_2D_depth =
@@ -192,24 +129,12 @@ static const char *get_appropriate_write_kernel(GLenum target,
 {
     switch (get_base_gl_target(target))
     {
-        case GL_TEXTURE_1D:
-
-            if (type == kHalf)
-                return kernelpattern_image_write_1D_half;
-            else
-                return kernelpattern_image_write_1D;
-            break;
+        case GL_TEXTURE_1D: return kernelpattern_image_write_1D; break;
         case GL_TEXTURE_BUFFER:
-            if (type == kHalf)
-                return kernelpattern_image_write_1D_buffer_half;
-            else
-                return kernelpattern_image_write_1D_buffer;
+            return kernelpattern_image_write_1D_buffer;
             break;
         case GL_TEXTURE_1D_ARRAY:
-            if (type == kHalf)
-                return kernelpattern_image_write_1Darray_half;
-            else
-                return kernelpattern_image_write_1Darray;
+            return kernelpattern_image_write_1Darray;
             break;
         case GL_COLOR_ATTACHMENT0:
         case GL_RENDERBUFFER:
@@ -220,29 +145,18 @@ static const char *get_appropriate_write_kernel(GLenum target,
             if (channel_order == CL_DEPTH || channel_order == CL_DEPTH_STENCIL)
                 return kernelpattern_image_write_2D_depth;
 #endif
-            if (type == kHalf)
-                return kernelpattern_image_write_2D_half;
-            else
-                return kernelpattern_image_write_2D;
+            return kernelpattern_image_write_2D;
             break;
 
         case GL_TEXTURE_2D_ARRAY:
 #ifdef GL_VERSION_3_2
             if (channel_order == CL_DEPTH || channel_order == CL_DEPTH_STENCIL)
                 return kernelpattern_image_write_2D_array_depth;
 #endif
-            if (type == kHalf)
-                return kernelpattern_image_write_2Darray_half;
-            else
-                return kernelpattern_image_write_2Darray;
+            return kernelpattern_image_write_2Darray;
             break;
 
-        case GL_TEXTURE_3D:
-            if (type == kHalf)
-                return kernelpattern_image_write_3D_half;
-            else
-                return kernelpattern_image_write_3D;
-            break;
+        case GL_TEXTURE_3D: return kernelpattern_image_write_3D; break;
 
         default:
             log_error("Unsupported GL tex target (%s) passed to write test: "
@@ -310,8 +224,7 @@ void set_dimensions_by_target(GLenum target, size_t *dims, size_t sizes[3],
 int test_cl_image_write(cl_context context, cl_command_queue queue,
                         GLenum target, cl_mem clImage, size_t width,
                         size_t height, size_t depth, cl_image_format *outFormat,
-                        ExplicitType *outType, void **outSourceBuffer, MTdata d,
-                        bool supports_half)
+                        ExplicitType *outType, void **outSourceBuffer, MTdata d)
 {
     size_t global_dims, global_sizes[3];
     clProgramWrapper program;
@@ -335,11 +248,6 @@ int test_cl_image_write(cl_context context, cl_command_queue queue,
 
     const char *appropriateKernel = get_appropriate_write_kernel(
         target, *outType, outFormat->image_channel_order);
-    if (*outType == kHalf && !supports_half)
-    {
-        log_info("cl_khr_fp16 isn't supported. Skip this test.\n");
-        return 0;
-    }
 
     const char *suffix = get_kernel_suffix(outFormat);
     const char *convert = get_write_conversion(outFormat, *outType);
@@ -429,8 +337,7 @@ static int test_image_write(cl_context context, cl_command_queue queue,
                             GLenum glTarget, GLuint glTexture, size_t width,
                             size_t height, size_t depth,
                             cl_image_format *outFormat, ExplicitType *outType,
-                            void **outSourceBuffer, MTdata d,
-                            bool supports_half)
+                            void **outSourceBuffer, MTdata d)
 {
     int error;
 
@@ -450,8 +357,7 @@ static int test_image_write(cl_context context, cl_command_queue queue,
     }
 
     return test_cl_image_write(context, queue, glTarget, image, width, height,
-                               depth, outFormat, outType, outSourceBuffer, d,
-                               supports_half);
+                               depth, outFormat, outType, outSourceBuffer, d);
 }
 
 int supportsHalf(cl_context context, bool *supports_half)
@@ -523,11 +429,6 @@ static int test_image_format_write(cl_context context, cl_command_queue queue,
                                    ExplicitType type, MTdata d)
 {
     int error;
-    // If we're testing a half float format, then we need to determine the
-    // rounding mode of this machine.  Punt if we fail to do so.
-
-    if (type == kHalf)
-        if (DetectFloatToHalfRoundingMode(queue)) return 1;
 
     // Create an appropriate GL texture or renderbuffer, given the target.
 
@@ -616,28 +517,19 @@ static int test_image_format_write(cl_context context, cl_command_queue queue,
         globj = glRenderbuffer;
     }
 
-    bool supports_half = false;
-    error = supportsHalf(context, &supports_half);
-    if (error != 0) return error;
-
-    error = test_image_write(context, queue, target, globj, width, height,
-                             depth, &clFormat, &sourceType,
-                             (void **)&outSourceBuffer, d, supports_half);
+    error =
+        test_image_write(context, queue, target, globj, width, height, depth,
+                         &clFormat, &sourceType, (void **)&outSourceBuffer, d);
 
-    if (error != 0 || ((sourceType == kHalf) && !supports_half))
+    if (error != 0)
     {
         if (outSourceBuffer) free(outSourceBuffer);
         return error;
     }
 
     if (!outSourceBuffer) return 0;
 
-    // If actual source type was half, convert to float for validation.
-
-    if (sourceType == kHalf)
-        validationType = kFloat;
-    else
-        validationType = sourceType;
+    validationType = sourceType;
 
     BufferOwningPtr<char> validationSource;
 
@@ -685,7 +577,7 @@ static int test_image_format_write(cl_context context, cl_command_queue queue,
     int valid = 0;
     if (convertedGLResults)
     {
-        if (sourceType == kFloat || sourceType == kHalf)
+        if (sourceType == kFloat)
         {
             if (clFormat.image_channel_data_type == CL_UNORM_INT_101010)
             {

test_conformance/gl/test_renderbuffer.cpp

@@ -55,7 +55,7 @@ extern int test_cl_image_write(cl_context context, cl_command_queue queue,
                                size_t height, size_t depth,
                                cl_image_format *outFormat,
                                ExplicitType *outType, void **outSourceBuffer,
-                               MTdata d, bool supports_half);
+                               MTdata d);
 
 extern int test_cl_image_read(cl_context context, cl_command_queue queue,
                               GLenum gl_target, cl_mem image, size_t width,
@@ -299,7 +299,7 @@ int test_attach_renderbuffer_write_to_image(
     cl_context context, cl_command_queue queue, GLenum glTarget,
     GLuint glRenderbuffer, size_t imageWidth, size_t imageHeight,
     cl_image_format *outFormat, ExplicitType *outType, MTdata d,
-    void **outSourceBuffer, bool supports_half)
+    void **outSourceBuffer)
 {
     int error;
 
@@ -314,7 +314,7 @@ int test_attach_renderbuffer_write_to_image(
 
     return test_cl_image_write(context, queue, glTarget, image, imageWidth,
                                imageHeight, 1, outFormat, outType,
-                               outSourceBuffer, d, supports_half);
+                               outSourceBuffer, d);
 }
 
 int test_renderbuffer_image_write(cl_context context, cl_command_queue queue,
@@ -325,9 +325,6 @@ int test_renderbuffer_image_write(cl_context context, cl_command_queue queue,
 {
     int error;
 
-    if (type == kHalf)
-        if (DetectFloatToHalfRoundingMode(queue)) return 1;
-
     // Create the GL renderbuffer
     glFramebufferWrapper glFramebuffer;
     glRenderbufferWrapper glRenderbuffer;
@@ -355,20 +352,12 @@ int test_renderbuffer_image_write(cl_context context, cl_command_queue queue,
     ExplicitType validationType;
     void *outSourceBuffer;
 
-    bool supports_half = false;
-    error = supportsHalf(context, &supports_half);
-    if (error != 0) return error;
-
     error = test_attach_renderbuffer_write_to_image(
         context, queue, attachment, glRenderbuffer, width, height, &clFormat,
-        &sourceType, d, (void **)&outSourceBuffer, supports_half);
-    if (error != 0 || ((sourceType == kHalf) && !supports_half)) return error;
+        &sourceType, d, (void **)&outSourceBuffer);
+    if (error != 0) return error;
 
-    // If actual source type was half, convert to float for validation.
-    if (sourceType == kHalf)
-        validationType = kFloat;
-    else
-        validationType = sourceType;
+    validationType = sourceType;
 
     BufferOwningPtr<char> validationSource(convert_to_expected(
         outSourceBuffer, width * height, sourceType, validationType,
@@ -461,7 +450,6 @@ int test_renderbuffer_write(cl_device_id device, cl_context context,
         { GL_RGBA32UI_EXT, GL_RGBA_INTEGER_EXT, GL_UNSIGNED_INT, kUInt },
 #endif
         { GL_RGBA32F_ARB, GL_RGBA, GL_FLOAT, kFloat },
-        { GL_RGBA16F_ARB, GL_RGBA, GL_HALF_FLOAT, kHalf }
     };
 
     size_t fmtIdx, attIdx;