Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Removed usage of half types in CTS test #2079

Open
wants to merge 1 commit into
base: main
Choose a base branch
from
Open

Removed usage of half types in CTS test #2079

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 0 additions & 1 deletion test_conformance/gl/common.h
View file
Open in desktop
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
Open in desktop
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
Open in desktop
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
Loading