ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-19 06:09:01 +00:00
Code Issues Packages Projects Releases Wiki Activity

Removed usage of half types in CTS gl tests (#2147)

Browse Source 
see #1982
This commit is contained in:
Kamil-Goras-Mobica
2024-11-26 17:39:42 +01:00
committed by GitHub
parent e9a248f555
commit feca4c6354
2 changed files with 24 additions and 46 deletions
Download Patch File Download Diff File Expand all files Collapse all files

56
test_conformance/gl/test_images_write_common.cpp
View File

@@ -37,11 +37,10 @@ static const char *kernelpattern_image_write_1D =
"}\n"; "}\n";
static const char *kernelpattern_image_write_1D_half = static const char *kernelpattern_image_write_1D_half =
"#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n" "__kernel void sample_test( __global half *source, write_only image1d_t dest )\n"
"__kernel void sample_test( __global half4 *source, write_only image1d_t dest )\n"
"{\n" "{\n"
" uint index = get_global_id(0);\n" " uint index = get_global_id(0);\n"
" write_imagef( dest, index, vload_half4(index, (__global half *)source));\n" " write_imagef( dest, index, vload_half4(index, source));\n"
"}\n"; "}\n";
static const char *kernelpattern_image_write_1D_buffer = static const char *kernelpattern_image_write_1D_buffer =
@@ -53,11 +52,10 @@ static const char *kernelpattern_image_write_1D_buffer =
"}\n"; "}\n";
static const char *kernelpattern_image_write_1D_buffer_half = static const char *kernelpattern_image_write_1D_buffer_half =
"#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n" "__kernel void sample_test( __global half *source, write_only image1d_buffer_t dest )\n"
"__kernel void sample_test( __global half4 *source, write_only image1d_buffer_t dest )\n"
"{\n" "{\n"
" uint index = get_global_id(0);\n" " uint index = get_global_id(0);\n"
" write_imagef( dest, index, vload_half4(index, (__global half *)source));\n" " write_imagef( dest, index, vload_half4(index, source));\n"
"}\n"; "}\n";
static const char *kernelpattern_image_write_2D = static const char *kernelpattern_image_write_2D =
@@ -71,13 +69,12 @@ static const char *kernelpattern_image_write_2D =
"}\n"; "}\n";
static const char *kernelpattern_image_write_2D_half = static const char *kernelpattern_image_write_2D_half =
"#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n" "__kernel void sample_test( __global half *source, write_only image2d_t dest )\n"
"__kernel void sample_test( __global half4 *source, write_only image2d_t dest )\n"
"{\n" "{\n"
" int tidX = get_global_id(0);\n" " int tidX = get_global_id(0);\n"
" int tidY = get_global_id(1);\n" " int tidY = get_global_id(1);\n"
" uint index = tidY * get_image_width( dest ) + tidX;\n" " uint index = tidY * get_image_width( dest ) + tidX;\n"
" write_imagef( dest, (int2)( tidX, tidY ), vload_half4(index, (__global half *)source));\n" " write_imagef( dest, (int2)( tidX, tidY ), vload_half4(index, source));\n"
"}\n"; "}\n";
static const char *kernelpattern_image_write_1Darray = static const char *kernelpattern_image_write_1Darray =
@@ -91,13 +88,12 @@ static const char *kernelpattern_image_write_1Darray =
"}\n"; "}\n";
static const char *kernelpattern_image_write_1Darray_half = static const char *kernelpattern_image_write_1Darray_half =
"#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n" "__kernel void sample_test( __global half *source, write_only image1d_array_t dest )\n"
"__kernel void sample_test( __global half4 *source, write_only image1d_array_t dest )\n"
"{\n" "{\n"
" int tidX = get_global_id(0);\n" " int tidX = get_global_id(0);\n"
" int tidY = get_global_id(1);\n" " int tidY = get_global_id(1);\n"
" uint index = tidY * get_image_width( dest ) + tidX;\n" " uint index = tidY * get_image_width( dest ) + tidX;\n"
" write_imagef( dest, (int2)( tidX, tidY ), vload_half4(index, (__global half *)source));\n" " write_imagef( dest, (int2)( tidX, tidY ), vload_half4(index, source));\n"
"}\n"; "}\n";
static const char *kernelpattern_image_write_3D = static const char *kernelpattern_image_write_3D =
@@ -115,9 +111,8 @@ static const char *kernelpattern_image_write_3D =
"}\n"; "}\n";
static const char *kernelpattern_image_write_3D_half = 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" "#pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable\n"
"__kernel void sample_test( __global half4 *source, write_only image3d_t dest )\n" "__kernel void sample_test( __global half *source, write_only image3d_t dest )\n"
"{\n" "{\n"
" int tidX = get_global_id(0);\n" " int tidX = get_global_id(0);\n"
" int tidY = get_global_id(1);\n" " int tidY = get_global_id(1);\n"
@@ -125,7 +120,7 @@ static const char *kernelpattern_image_write_3D_half =
" int width = get_image_width( dest );\n" " int width = get_image_width( dest );\n"
" int height = get_image_height( dest );\n" " int height = get_image_height( dest );\n"
" int index = tidZ * width * height + tidY * width + tidX;\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" " write_imagef( dest, (int4)( tidX, tidY, tidZ, 0 ), vload_half4(index, source));\n"
"}\n"; "}\n";
static const char *kernelpattern_image_write_2Darray = static const char *kernelpattern_image_write_2Darray =
@@ -142,8 +137,7 @@ static const char *kernelpattern_image_write_2Darray =
"}\n"; "}\n";
static const char *kernelpattern_image_write_2Darray_half = static const char *kernelpattern_image_write_2Darray_half =
"#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n" "__kernel void sample_test( __global half *source, write_only image2d_array_t dest )\n"
"__kernel void sample_test( __global half4 *source, write_only image2d_array_t dest )\n"
"{\n" "{\n"
" int tidX = get_global_id(0);\n" " int tidX = get_global_id(0);\n"
" int tidY = get_global_id(1);\n" " int tidY = get_global_id(1);\n"
@@ -151,7 +145,7 @@ static const char *kernelpattern_image_write_2Darray_half =
" int width = get_image_width( dest );\n" " int width = get_image_width( dest );\n"
" int height = get_image_height( dest );\n" " int height = get_image_height( dest );\n"
" int index = tidZ * width * height + tidY * width + tidX;\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" " write_imagef( dest, (int4)( tidX, tidY, tidZ, 0 ), vload_half4(index, source));\n"
"}\n"; "}\n";
#ifdef GL_VERSION_3_2 #ifdef GL_VERSION_3_2
@@ -310,8 +304,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, int test_cl_image_write(cl_context context, cl_command_queue queue,
GLenum target, cl_mem clImage, size_t width, GLenum target, cl_mem clImage, size_t width,
size_t height, size_t depth, cl_image_format *outFormat, size_t height, size_t depth, cl_image_format *outFormat,
ExplicitType *outType, void **outSourceBuffer, MTdata d, ExplicitType *outType, void **outSourceBuffer, MTdata d)
bool supports_half)
{ {
size_t global_dims, global_sizes[3]; size_t global_dims, global_sizes[3];
clProgramWrapper program; clProgramWrapper program;
@@ -335,11 +328,6 @@ int test_cl_image_write(cl_context context, cl_command_queue queue,
const char *appropriateKernel = get_appropriate_write_kernel( const char *appropriateKernel = get_appropriate_write_kernel(
target, *outType, outFormat->image_channel_order); 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 *suffix = get_kernel_suffix(outFormat);
const char *convert = get_write_conversion(outFormat, *outType); const char *convert = get_write_conversion(outFormat, *outType);
@@ -429,8 +417,7 @@ static int test_image_write(cl_context context, cl_command_queue queue,
GLenum glTarget, GLuint glTexture, size_t width, GLenum glTarget, GLuint glTexture, size_t width,
size_t height, size_t depth, size_t height, size_t depth,
cl_image_format *outFormat, ExplicitType *outType, cl_image_format *outFormat, ExplicitType *outType,
void **outSourceBuffer, MTdata d, void **outSourceBuffer, MTdata d)
bool supports_half)
{ {
int error; int error;
@@ -450,8 +437,7 @@ static int test_image_write(cl_context context, cl_command_queue queue,
} }
return test_cl_image_write(context, queue, glTarget, image, width, height, return test_cl_image_write(context, queue, glTarget, image, width, height,
depth, outFormat, outType, outSourceBuffer, d, depth, outFormat, outType, outSourceBuffer, d);
supports_half);
} }
int supportsHalf(cl_context context, bool *supports_half) int supportsHalf(cl_context context, bool *supports_half)
@@ -616,15 +602,11 @@ static int test_image_format_write(cl_context context, cl_command_queue queue,
globj = glRenderbuffer; globj = glRenderbuffer;
} }
bool supports_half = false; error =
error = supportsHalf(context, &supports_half); test_image_write(context, queue, target, globj, width, height, depth,
if (error != 0) return error; &clFormat, &sourceType, (void **)&outSourceBuffer, d);
error = test_image_write(context, queue, target, globj, width, height, if (error != 0)
depth, &clFormat, &sourceType,
(void **)&outSourceBuffer, d, supports_half);
if (error != 0 || ((sourceType == kHalf) && !supports_half))
{ {
if (outSourceBuffer) free(outSourceBuffer); if (outSourceBuffer) free(outSourceBuffer);
return error; return error;

14
test_conformance/gl/test_renderbuffer.cpp
View File

@@ -55,7 +55,7 @@ extern int test_cl_image_write(cl_context context, cl_command_queue queue,
size_t height, size_t depth, size_t height, size_t depth,
cl_image_format *outFormat, cl_image_format *outFormat,
ExplicitType *outType, void **outSourceBuffer, ExplicitType *outType, void **outSourceBuffer,
MTdata d, bool supports_half); MTdata d);
extern int test_cl_image_read(cl_context context, cl_command_queue queue, extern int test_cl_image_read(cl_context context, cl_command_queue queue,
GLenum gl_target, cl_mem image, size_t width, 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, cl_context context, cl_command_queue queue, GLenum glTarget,
GLuint glRenderbuffer, size_t imageWidth, size_t imageHeight, GLuint glRenderbuffer, size_t imageWidth, size_t imageHeight,
cl_image_format *outFormat, ExplicitType *outType, MTdata d, cl_image_format *outFormat, ExplicitType *outType, MTdata d,
void **outSourceBuffer, bool supports_half) void **outSourceBuffer)
{ {
int error; int error;
@@ -314,7 +314,7 @@ int test_attach_renderbuffer_write_to_image(
return test_cl_image_write(context, queue, glTarget, image, imageWidth, return test_cl_image_write(context, queue, glTarget, image, imageWidth,
imageHeight, 1, outFormat, outType, imageHeight, 1, outFormat, outType,
outSourceBuffer, d, supports_half); outSourceBuffer, d);
} }
int test_renderbuffer_image_write(cl_context context, cl_command_queue queue, int test_renderbuffer_image_write(cl_context context, cl_command_queue queue,
@@ -355,14 +355,10 @@ int test_renderbuffer_image_write(cl_context context, cl_command_queue queue,
ExplicitType validationType; ExplicitType validationType;
void *outSourceBuffer; void *outSourceBuffer;
bool supports_half = false;
error = supportsHalf(context, &supports_half);
if (error != 0) return error;
error = test_attach_renderbuffer_write_to_image( error = test_attach_renderbuffer_write_to_image(
context, queue, attachment, glRenderbuffer, width, height, &clFormat, context, queue, attachment, glRenderbuffer, width, height, &clFormat,
&sourceType, d, (void **)&outSourceBuffer, supports_half); &sourceType, d, (void **)&outSourceBuffer);
if (error != 0 || ((sourceType == kHalf) && !supports_half)) return error; if (error != 0) return error;
// If actual source type was half, convert to float for validation. // If actual source type was half, convert to float for validation.
if (sourceType == kHalf) if (sourceType == kHalf)