Skip to content

Commit

Permalink
Removed usage of half types in CTS test #1982
Browse files Browse the repository at this point in the history
  • Loading branch information
@kamil-goras-mobica
kamil-goras-mobica committed Sep 11, 2024
1 parent 0a00a1f commit 21218ff
Show file tree
Hide file tree
Showing 3 changed files with 21 additions and 142 deletions.
1 change: 0 additions & 1 deletion test_conformance/gl/common.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
138 changes: 15 additions & 123 deletions test_conformance/gl/test_images_write_common.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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"
Expand All @@ -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"
Expand All @@ -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"
Expand All @@ -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"
Expand All @@ -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"
Expand All @@ -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 =
Expand Down Expand Up @@ -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:
Expand All @@ -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: "
Expand Down Expand Up @@ -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;
Expand All @@ -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);
Expand Down Expand Up @@ -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;

Expand All @@ -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)
Expand Down Expand Up @@ -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.

Expand Down Expand Up @@ -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;

Expand Down Expand Up @@ -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)
{
Expand Down
24 changes: 6 additions & 18 deletions test_conformance/gl/test_renderbuffer.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down Expand Up @@ -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;

Expand All @@ -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,
Expand All @@ -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;
Expand Down Expand Up @@ -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,
Expand Down Expand Up @@ -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;
Expand Down

0 comments on commit 21218ff

Please sign in to comment.