aboutsummaryrefslogtreecommitdiff
path: root/test_conformance/select/test_select.cpp
diff options
context:
space:
mode:
Diffstat (limited to 'test_conformance/select/test_select.cpp')
-rw-r--r--test_conformance/select/test_select.cpp320
1 files changed, 163 insertions, 157 deletions
diff --git a/test_conformance/select/test_select.cpp b/test_conformance/select/test_select.cpp
index 7fa3bc08..72be08c7 100644
--- a/test_conformance/select/test_select.cpp
+++ b/test_conformance/select/test_select.cpp
@@ -14,11 +14,16 @@
// limitations under the License.
//
#include "harness/compat.h"
+#include "harness/typeWrappers.h"
#include <assert.h>
#include <stdio.h>
#include <time.h>
#include <string.h>
+
+#include <cinttypes>
+#include <vector>
+
#if ! defined( _WIN32)
#if defined(__APPLE__)
#include <sys/sysctl.h>
@@ -42,11 +47,14 @@ static void initSrcBuffer(void* src1, Type stype, MTdata);
// initialize the valued used to compare with in the select with
// vlaues [start, count)
-static void initCmpBuffer(void* cmp, Type cmptype, uint64_t start, size_t count);
+static void initCmpBuffer(void *cmp, Type cmptype, uint64_t start,
+ const size_t count);
// make a program that uses select for the given stype (src/dest type),
// ctype (comparison type), veclen (vector length)
-static cl_program makeSelectProgram(cl_kernel *kernel_ptr, const cl_context context, Type stype, Type ctype, size_t veclen );
+static cl_program makeSelectProgram(cl_kernel *kernel_ptr, cl_context context,
+ Type stype, Type ctype,
+ const size_t veclen);
// Creates and execute the select test for the given device, context,
// stype (source/dest type), cmptype (comparison type), using max_tg_size
@@ -66,6 +74,16 @@ static void printUsage( void );
#define BUFFER_SIZE (1024*1024)
#define KPAGESIZE 4096
+#define test_error_count(errCode, msg) \
+ { \
+ auto errCodeResult = errCode; \
+ if (errCodeResult != CL_SUCCESS) \
+ { \
+ gFailCount++; \
+ print_error(errCodeResult, msg); \
+ return errCode; \
+ } \
+ }
// When we indicate non wimpy mode, the types that are 32 bits value will
// test their entire range and 64 bits test will test the 32 bit
@@ -74,12 +92,6 @@ static void printUsage( void );
static bool s_wimpy_mode = false;
static int s_wimpy_reduction_factor = 256;
-// Tests are broken into the major test which is based on the
-// src and cmp type and their corresponding vector types and
-// sub tests which is for each individual test. The following
-// tracks the subtests
-int s_test_cnt = 0;
-
//-----------------------------------------
// Static helper functions
//-----------------------------------------
@@ -112,36 +124,37 @@ static void initSrcBuffer(void* src1, Type stype, MTdata d)
s1[i] = genrand_int32(d);
}
-static void initCmpBuffer(void* cmp, Type cmptype, uint64_t start, size_t count) {
- int i;
+static void initCmpBuffer(void *cmp, Type cmptype, uint64_t start,
+ const size_t count)
+
+{
assert(cmptype != kfloat);
switch (type_size[cmptype]) {
case 1: {
uint8_t* ub = (uint8_t *)cmp;
- for (i=0; i < count; ++i)
- ub[i] = (uint8_t)start++;
+ for (size_t i = 0; i < count; ++i) ub[i] = (uint8_t)start++;
break;
}
case 2: {
uint16_t* us = (uint16_t *)cmp;
- for (i=0; i < count; ++i)
- us[i] = (uint16_t)start++;
+ for (size_t i = 0; i < count; ++i) us[i] = (uint16_t)start++;
break;
}
case 4: {
if (!s_wimpy_mode) {
uint32_t* ui = (uint32_t *)cmp;
- for (i=0; i < count; ++i)
- ui[i] = (uint32_t)start++;
+ for (size_t i = 0; i < count; ++i) ui[i] = (uint32_t)start++;
}
else {
// The short test doesn't iterate over the entire 32 bit space so
// we alternate between positive and negative values
int32_t* ui = (int32_t *)cmp;
- int32_t sign = 1;
- for (i=0; i < count; ++i, ++start) {
- ui[i] = (int32_t)start*sign;
- sign = sign * -1;
+ int32_t neg_start = (int32_t)start * -1;
+ for (size_t i = 0; i < count; i++)
+ {
+ ++start;
+ --neg_start;
+ ui[i] = (int32_t)((i % 2) ? start : neg_start);
}
}
break;
@@ -150,10 +163,12 @@ static void initCmpBuffer(void* cmp, Type cmptype, uint64_t start, size_t count)
// We don't iterate over the entire space of 64 bit so for the
// selects, we want to test positive and negative values
int64_t* ll = (int64_t *)cmp;
- int64_t sign = 1;
- for (i=0; i < count; ++i, ++start) {
- ll[i] = start*sign;
- sign = sign * -1;
+ int64_t neg_start = (int64_t)start * -1;
+ for (size_t i = 0; i < count; i++)
+ {
+ ++start;
+ --neg_start;
+ ll[i] = (int64_t)((i % 2) ? start : neg_start);
}
break;
}
@@ -165,7 +180,9 @@ static void initCmpBuffer(void* cmp, Type cmptype, uint64_t start, size_t count)
// Make the various incarnations of the program we want to run
// stype: source and destination type for the select
// ctype: compare type
-static cl_program makeSelectProgram(cl_kernel *kernel_ptr, const cl_context context, Type srctype, Type cmptype, size_t vec_len)
+static cl_program makeSelectProgram(cl_kernel *kernel_ptr,
+ const cl_context context, Type srctype,
+ Type cmptype, const size_t vec_len)
{
char testname[256];
char stypename[32];
@@ -237,6 +254,9 @@ static cl_program makeSelectProgram(cl_kernel *kernel_ptr, const cl_context cont
if (srctype == kdouble)
strcpy( extension, "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n" );
+ if (srctype == khalf)
+ strcpy(extension, "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n");
+
// create type name and testname
switch( vec_len )
{
@@ -288,39 +308,38 @@ static cl_program makeSelectProgram(cl_kernel *kernel_ptr, const cl_context cont
return program;
}
-
#define VECTOR_SIZE_COUNT 6
static int doTest(cl_command_queue queue, cl_context context, Type stype, Type cmptype, cl_device_id device)
{
int err = CL_SUCCESS;
- int s_test_fail = 0;
- MTdataHolder d;
+ MTdataHolder d(gRandomSeed);
const size_t element_count[VECTOR_SIZE_COUNT] = { 1, 2, 3, 4, 8, 16 };
- cl_mem src1 = NULL;
- cl_mem src2 = NULL;
- cl_mem cmp = NULL;
- cl_mem dest = NULL;
- void *ref = NULL;
- void *sref = NULL;
+ clMemWrapper src1, src2, cmp, dest;
cl_ulong blocks = type_size[stype] * 0x100000000ULL / BUFFER_SIZE;
- size_t block_elements = BUFFER_SIZE / type_size[stype];
+ const size_t block_elements = BUFFER_SIZE / type_size[stype];
size_t step = s_wimpy_mode ? s_wimpy_reduction_factor : 1;
cl_ulong cmp_stride = block_elements * step;
// It is more efficient to create the tests all at once since we
// use the same test data on each of the vector sizes
- int vecsize;
- cl_program programs[VECTOR_SIZE_COUNT];
- cl_kernel kernels[VECTOR_SIZE_COUNT];
+ clProgramWrapper programs[VECTOR_SIZE_COUNT];
+ clKernelWrapper kernels[VECTOR_SIZE_COUNT];
- if(stype == kdouble && ! is_extension_available( device, "cl_khr_fp64" ))
+ if (stype == kdouble && !is_extension_available(device, "cl_khr_fp64"))
{
log_info("Skipping double because cl_khr_fp64 extension is not supported.\n");
return 0;
}
+ if (stype == khalf && !is_extension_available(device, "cl_khr_fp16"))
+ {
+ log_info(
+ "Skipping half because cl_khr_fp16 extension is not supported.\n");
+ return 0;
+ }
+
if (gIsEmbedded)
{
if (( stype == klong || stype == kulong ) && ! is_extension_available( device, "cles_khr_int64" ))
@@ -336,29 +355,51 @@ static int doTest(cl_command_queue queue, cl_context context, Type stype, Type c
}
}
- for (vecsize = 0; vecsize < VECTOR_SIZE_COUNT; ++vecsize)
- {
- programs[vecsize] = makeSelectProgram(&kernels[vecsize], context, stype, cmptype, element_count[vecsize] );
- if (!programs[vecsize] || !kernels[vecsize]) {
- ++s_test_fail;
- ++s_test_cnt;
- return -1;
- }
- }
-
- ref = malloc( BUFFER_SIZE );
- if( NULL == ref ){ log_error("Error: could not allocate ref buffer\n" ); goto exit; }
- sref = malloc( BUFFER_SIZE );
- if( NULL == sref ){ log_error("Error: could not allocate ref buffer\n" ); goto exit; }
src1 = clCreateBuffer( context, CL_MEM_READ_ONLY, BUFFER_SIZE, NULL, &err );
- if( err ) { log_error( "Error: could not allocate src1 buffer\n" ); ++s_test_fail; goto exit; }
+ test_error_count(err, "Error: could not allocate src1 buffer\n");
src2 = clCreateBuffer( context, CL_MEM_READ_ONLY, BUFFER_SIZE, NULL, &err );
- if( err ) { log_error( "Error: could not allocate src2 buffer\n" ); ++s_test_fail; goto exit; }
+ test_error_count(err, "Error: could not allocate src2 buffer\n");
cmp = clCreateBuffer( context, CL_MEM_READ_ONLY, BUFFER_SIZE, NULL, &err );
- if( err ) { log_error( "Error: could not allocate cmp buffer\n" ); ++s_test_fail; goto exit; }
+ test_error_count(err, "Error: could not allocate cmp buffer\n");
dest = clCreateBuffer( context, CL_MEM_WRITE_ONLY, BUFFER_SIZE, NULL, &err );
- if( err ) { log_error( "Error: could not allocate dest buffer\n" ); ++s_test_fail; goto exit; }
+ test_error_count(err, "Error: could not allocate dest buffer\n");
+
+ programs[0] = makeSelectProgram(&kernels[0], context, stype, cmptype,
+ element_count[0]);
+ programs[1] = makeSelectProgram(&kernels[1], context, stype, cmptype,
+ element_count[1]);
+ programs[2] = makeSelectProgram(&kernels[2], context, stype, cmptype,
+ element_count[2]);
+ programs[3] = makeSelectProgram(&kernels[3], context, stype, cmptype,
+ element_count[3]);
+ programs[4] = makeSelectProgram(&kernels[4], context, stype, cmptype,
+ element_count[4]);
+ programs[5] = makeSelectProgram(&kernels[5], context, stype, cmptype,
+ element_count[5]);
+
+ for (size_t vecsize = 0; vecsize < VECTOR_SIZE_COUNT; ++vecsize)
+ {
+ if (!programs[vecsize] || !kernels[vecsize])
+ {
+ return -1;
+ }
+
+ err = clSetKernelArg(kernels[vecsize], 0, sizeof dest, &dest);
+ test_error_count(err, "Error: Cannot set kernel arg dest!\n");
+ err = clSetKernelArg(kernels[vecsize], 1, sizeof src1, &src1);
+ test_error_count(err, "Error: Cannot set kernel arg dest!\n");
+ err = clSetKernelArg(kernels[vecsize], 2, sizeof src2, &src2);
+ test_error_count(err, "Error: Cannot set kernel arg dest!\n");
+ err = clSetKernelArg(kernels[vecsize], 3, sizeof cmp, &cmp);
+ test_error_count(err, "Error: Cannot set kernel arg dest!\n");
+ }
+ std::vector<char> ref(BUFFER_SIZE);
+ std::vector<char> sref(BUFFER_SIZE);
+ std::vector<char> src1_host(BUFFER_SIZE);
+ std::vector<char> src2_host(BUFFER_SIZE);
+ std::vector<char> cmp_host(BUFFER_SIZE);
+ std::vector<char> dest_host(BUFFER_SIZE);
// We block the test as we are running over the range of compare values
// "block the test" means "break the test into blocks"
@@ -368,81 +409,63 @@ static int doTest(cl_command_queue queue, cl_context context, Type stype, Type c
cmp_stride = block_elements * step * (0xffffffffffffffffULL / 0x100000000ULL + 1);
log_info("Testing...");
- d = MTdataHolder(gRandomSeed);
uint64_t i;
+
+ initSrcBuffer(src1_host.data(), stype, d);
+ initSrcBuffer(src2_host.data(), stype, d);
for (i=0; i < blocks; i+=step)
{
- void *s1 = clEnqueueMapBuffer( queue, src1, CL_TRUE, CL_MAP_WRITE, 0, BUFFER_SIZE, 0, NULL, NULL, &err );
- if( err ){ log_error( "Error: Could not map src1" ); goto exit; }
- // Setup the input data to change for each block
- initSrcBuffer( s1, stype, d);
-
- void *s2 = clEnqueueMapBuffer( queue, src2, CL_TRUE, CL_MAP_WRITE, 0, BUFFER_SIZE, 0, NULL, NULL, &err );
- if( err ){ log_error( "Error: Could not map src2" ); goto exit; }
- // Setup the input data to change for each block
- initSrcBuffer( s2, stype, d);
-
- void *s3 = clEnqueueMapBuffer( queue, cmp, CL_TRUE, CL_MAP_WRITE, 0, BUFFER_SIZE, 0, NULL, NULL, &err );
- if( err ){ log_error( "Error: Could not map cmp" ); goto exit; }
- // Setup the input data to change for each block
- initCmpBuffer(s3, cmptype, i * cmp_stride, block_elements);
-
- // Create the reference result
- Select sfunc = (cmptype == ctype[stype][0]) ? vrefSelects[stype][0] : vrefSelects[stype][1];
- (*sfunc)(ref, s1, s2, s3, block_elements);
-
- sfunc = (cmptype == ctype[stype][0]) ? refSelects[stype][0] : refSelects[stype][1];
- (*sfunc)(sref, s1, s2, s3, block_elements);
-
- if( (err = clEnqueueUnmapMemObject( queue, src1, s1, 0, NULL, NULL )))
- { log_error( "Error: coult not unmap src1\n" ); ++s_test_fail; goto exit; }
- if( (err = clEnqueueUnmapMemObject( queue, src2, s2, 0, NULL, NULL )))
- { log_error( "Error: coult not unmap src2\n" ); ++s_test_fail; goto exit; }
- if( (err = clEnqueueUnmapMemObject( queue, cmp, s3, 0, NULL, NULL )))
- { log_error( "Error: coult not unmap cmp\n" ); ++s_test_fail; goto exit; }
-
- for (vecsize = 0; vecsize < VECTOR_SIZE_COUNT; ++vecsize)
+ initCmpBuffer(cmp_host.data(), cmptype, i * cmp_stride, block_elements);
+
+ err = clEnqueueWriteBuffer(queue, src1, CL_FALSE, 0, BUFFER_SIZE,
+ src1_host.data(), 0, NULL, NULL);
+ test_error_count(err, "Error: Could not write src1");
+
+ err = clEnqueueWriteBuffer(queue, src2, CL_FALSE, 0, BUFFER_SIZE,
+ src2_host.data(), 0, NULL, NULL);
+ test_error_count(err, "Error: Could not write src2");
+
+ err = clEnqueueWriteBuffer(queue, cmp, CL_FALSE, 0, BUFFER_SIZE,
+ cmp_host.data(), 0, NULL, NULL);
+ test_error_count(err, "Error: Could not write cmp");
+
+ Select sfunc = (cmptype == ctype[stype][0]) ? vrefSelects[stype][0]
+ : vrefSelects[stype][1];
+ (*sfunc)(ref.data(), src1_host.data(), src2_host.data(),
+ cmp_host.data(), block_elements);
+
+ sfunc = (cmptype == ctype[stype][0]) ? refSelects[stype][0]
+ : refSelects[stype][1];
+ (*sfunc)(sref.data(), src1_host.data(), src2_host.data(),
+ cmp_host.data(), block_elements);
+
+ for (int vecsize = 0; vecsize < VECTOR_SIZE_COUNT; ++vecsize)
{
size_t vector_size = element_count[vecsize] * type_size[stype];
size_t vector_count = (BUFFER_SIZE + vector_size - 1) / vector_size;
- if((err = clSetKernelArg(kernels[vecsize], 0, sizeof dest, &dest) ))
- { log_error( "Error: Cannot set kernel arg dest! %d\n", err ); ++s_test_fail; goto exit; }
- if((err = clSetKernelArg(kernels[vecsize], 1, sizeof src1, &src1) ))
- { log_error( "Error: Cannot set kernel arg dest! %d\n", err ); ++s_test_fail; goto exit; }
- if((err = clSetKernelArg(kernels[vecsize], 2, sizeof src2, &src2) ))
- { log_error( "Error: Cannot set kernel arg dest! %d\n", err ); ++s_test_fail; goto exit; }
- if((err = clSetKernelArg(kernels[vecsize], 3, sizeof cmp, &cmp) ))
- { log_error( "Error: Cannot set kernel arg dest! %d\n", err ); ++s_test_fail; goto exit; }
-
+ const cl_int pattern = -1;
+ err = clEnqueueFillBuffer(queue, dest, &pattern, sizeof(cl_int), 0,
+ BUFFER_SIZE, 0, nullptr, nullptr);
+ test_error_count(err, "clEnqueueFillBuffer failed");
- // Wipe destination
- void *d = clEnqueueMapBuffer( queue, dest, CL_TRUE, CL_MAP_WRITE, 0, BUFFER_SIZE, 0, NULL, NULL, &err );
- if( err ){ log_error( "Error: Could not map dest" ); ++s_test_fail; goto exit; }
- memset( d, -1, BUFFER_SIZE );
- if( (err = clEnqueueUnmapMemObject( queue, dest, d, 0, NULL, NULL ) ) ){ log_error( "Error: Could not unmap dest" ); ++s_test_fail; goto exit; }
err = clEnqueueNDRangeKernel(queue, kernels[vecsize], 1, NULL, &vector_count, NULL, 0, NULL, NULL);
- if (err != CL_SUCCESS) {
- log_error("clEnqueueNDRangeKernel failed errcode:%d\n", err);
- ++s_test_fail;
- goto exit;
- }
-
- d = clEnqueueMapBuffer( queue, dest, CL_TRUE, CL_MAP_READ, 0, BUFFER_SIZE, 0, NULL, NULL, &err );
- if( err ){ log_error( "Error: Could not map dest # 2" ); ++s_test_fail; goto exit; }
+ test_error_count(err, "clEnqueueNDRangeKernel failed errcode\n");
- if ((*checkResults[stype])(d, vecsize == 0 ? sref : ref, block_elements, element_count[vecsize])!=0){
- log_error("vec_size:%d indx: 0x%16.16llx\n", (int)element_count[vecsize], i);
- ++s_test_fail;
- goto exit;
- }
+ err = clEnqueueReadBuffer(queue, dest, CL_TRUE, 0, BUFFER_SIZE,
+ dest_host.data(), 0, NULL, NULL);
+ test_error_count(
+ err, "Error: Reading buffer from dest to dest_host failed\n");
- if( (err = clEnqueueUnmapMemObject( queue, dest, d, 0, NULL, NULL ) ) )
+ if ((*checkResults[stype])(dest_host.data(),
+ vecsize == 0 ? sref.data() : ref.data(),
+ block_elements, element_count[vecsize])
+ != 0)
{
- log_error( "Error: Could not unmap dest" );
- ++s_test_fail;
- goto exit;
+ log_error("vec_size:%d indx: 0x%16.16" PRIx64 "\n",
+ (int)element_count[vecsize], i);
+ return TEST_FAIL;
}
} // for vecsize
} // for i
@@ -452,24 +475,6 @@ static int doTest(cl_command_queue queue, cl_context context, Type stype, Type c
else
log_info(" Wimpy Passed\n\n");
-exit:
- if( src1 ) clReleaseMemObject( src1 );
- if( src2 ) clReleaseMemObject( src2 );
- if( cmp ) clReleaseMemObject( cmp );
- if( dest) clReleaseMemObject( dest );
- if( ref ) free(ref );
- if( sref ) free(sref );
-
- for (vecsize = 0; vecsize < VECTOR_SIZE_COUNT; vecsize++) {
- clReleaseKernel(kernels[vecsize]);
- clReleaseProgram(programs[vecsize]);
- }
- ++s_test_cnt;
- if (s_test_fail)
- {
- err = TEST_FAIL;
- gFailCount++;
- }
return err;
}
@@ -505,6 +510,16 @@ int test_select_short_short(cl_device_id deviceID, cl_context context, cl_comman
{
return doTest(queue, context, kshort, kshort, deviceID);
}
+int test_select_half_ushort(cl_device_id deviceID, cl_context context,
+ cl_command_queue queue, int num_elements)
+{
+ return doTest(queue, context, khalf, kushort, deviceID);
+}
+int test_select_half_short(cl_device_id deviceID, cl_context context,
+ cl_command_queue queue, int num_elements)
+{
+ return doTest(queue, context, khalf, kshort, deviceID);
+}
int test_select_uint_uint(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
{
return doTest(queue, context, kuint, kuint, deviceID);
@@ -555,26 +570,17 @@ int test_select_double_long(cl_device_id deviceID, cl_context context, cl_comman
}
test_definition test_list[] = {
- ADD_TEST( select_uchar_uchar ),
- ADD_TEST( select_uchar_char ),
- ADD_TEST( select_char_uchar ),
- ADD_TEST( select_char_char ),
- ADD_TEST( select_ushort_ushort ),
- ADD_TEST( select_ushort_short ),
- ADD_TEST( select_short_ushort ),
- ADD_TEST( select_short_short ),
- ADD_TEST( select_uint_uint ),
- ADD_TEST( select_uint_int ),
- ADD_TEST( select_int_uint ),
- ADD_TEST( select_int_int ),
- ADD_TEST( select_float_uint ),
- ADD_TEST( select_float_int ),
- ADD_TEST( select_ulong_ulong ),
- ADD_TEST( select_ulong_long ),
- ADD_TEST( select_long_ulong ),
- ADD_TEST( select_long_long ),
- ADD_TEST( select_double_ulong ),
- ADD_TEST( select_double_long ),
+ ADD_TEST(select_uchar_uchar), ADD_TEST(select_uchar_char),
+ ADD_TEST(select_char_uchar), ADD_TEST(select_char_char),
+ ADD_TEST(select_ushort_ushort), ADD_TEST(select_ushort_short),
+ ADD_TEST(select_short_ushort), ADD_TEST(select_short_short),
+ ADD_TEST(select_half_ushort), ADD_TEST(select_half_short),
+ ADD_TEST(select_uint_uint), ADD_TEST(select_uint_int),
+ ADD_TEST(select_int_uint), ADD_TEST(select_int_int),
+ ADD_TEST(select_float_uint), ADD_TEST(select_float_int),
+ ADD_TEST(select_ulong_ulong), ADD_TEST(select_ulong_long),
+ ADD_TEST(select_long_ulong), ADD_TEST(select_long_long),
+ ADD_TEST(select_double_ulong), ADD_TEST(select_double_long),
};
const int test_num = ARRAY_SIZE( test_list );
generated by cgit v1.2.3 (git 2.25.1) at 2024-06-09 16:37:47 +0000