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

remove min max macros (#1310)

Browse Source 
* remove the MIN and MAX macros and use the std versions instead

* fix formatting

* fix Arm build

* remove additional MIN and MAX macros from compat.h
This commit is contained in:
Ben Ashbaugh
2021-09-13 05:25:32 -07:00
committed by GitHub
parent 1f26e1d8ba
commit 02bf24d2b1
31 changed files with 241 additions and 202 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
test_common/harness/compat.h
View File

@@ -309,13 +309,6 @@ EXTERN_C int __builtin_clz(unsigned int pattern);
#endif
#ifndef MIN
#define MIN(x, y) (((x) < (y)) ? (x) : (y))
#endif
#ifndef MAX
#define MAX(x, y) (((x) > (y)) ? (x) : (y))
#endif
/*-----------------------------------------------------------------------------
WARNING: DO NOT USE THESE MACROS:

20
test_common/harness/errorHelpers.cpp
View File

@@ -18,6 +18,8 @@
#include <stdlib.h>
#include <string.h>
#include <algorithm>
#include "errorHelpers.h"
#include "parseParameters.h"
@@ -301,10 +303,6 @@ const char *GetQueuePropertyName(cl_command_queue_properties property)
}
}
#ifndef MAX
#define MAX(_a, _b) ((_a) > (_b) ? (_a) : (_b))
#endif
#if defined(_MSC_VER)
#define scalbnf(_a, _i) ldexpf(_a, _i)
#define scalbn(_a, _i) ldexp(_a, _i)
@@ -357,7 +355,7 @@ static float Ulp_Error_Half_Float(float test, double reference)
// The unbiased exponent of the ulp unit place
int ulp_exp =
HALF_MANT_DIG - 1 - MAX(ilogb(reference), HALF_MIN_EXP - 1);
HALF_MANT_DIG - 1 - std::max(ilogb(reference), HALF_MIN_EXP - 1);
// Scale the exponent of the error
return (float)scalbn(testVal - reference, ulp_exp);
@@ -365,7 +363,7 @@ static float Ulp_Error_Half_Float(float test, double reference)
// reference is a normal power of two or a zero
int ulp_exp =
HALF_MANT_DIG - 1 - MAX(ilogb(reference) - 1, HALF_MIN_EXP - 1);
HALF_MANT_DIG - 1 - std::max(ilogb(reference) - 1, HALF_MIN_EXP - 1);
// Scale the exponent of the error
return (float)scalbn(testVal - reference, ulp_exp);
@@ -437,7 +435,8 @@ float Ulp_Error(float test, double reference)
return 0.0f; // if we are expecting a NaN, any NaN is fine
// The unbiased exponent of the ulp unit place
int ulp_exp = FLT_MANT_DIG - 1 - MAX(ilogb(reference), FLT_MIN_EXP - 1);
int ulp_exp =
FLT_MANT_DIG - 1 - std::max(ilogb(reference), FLT_MIN_EXP - 1);
// Scale the exponent of the error
return (float)scalbn(testVal - reference, ulp_exp);
@@ -445,7 +444,8 @@ float Ulp_Error(float test, double reference)
// reference is a normal power of two or a zero
// The unbiased exponent of the ulp unit place
int ulp_exp = FLT_MANT_DIG - 1 - MAX(ilogb(reference) - 1, FLT_MIN_EXP - 1);
int ulp_exp =
FLT_MANT_DIG - 1 - std::max(ilogb(reference) - 1, FLT_MIN_EXP - 1);
// Scale the exponent of the error
return (float)scalbn(testVal - reference, ulp_exp);
@@ -513,7 +513,7 @@ float Ulp_Error_Double(double test, long double reference)
// The unbiased exponent of the ulp unit place
int ulp_exp =
DBL_MANT_DIG - 1 - MAX(ilogbl(reference), DBL_MIN_EXP - 1);
DBL_MANT_DIG - 1 - std::max(ilogbl(reference), DBL_MIN_EXP - 1);
// Scale the exponent of the error
float result = (float)scalbnl(testVal - reference, ulp_exp);
@@ -529,7 +529,7 @@ float Ulp_Error_Double(double test, long double reference)
// reference is a normal power of two or a zero
// The unbiased exponent of the ulp unit place
int ulp_exp =
DBL_MANT_DIG - 1 - MAX(ilogbl(reference) - 1, DBL_MIN_EXP - 1);
DBL_MANT_DIG - 1 - std::max(ilogbl(reference) - 1, DBL_MIN_EXP - 1);
// Scale the exponent of the error
float result = (float)scalbnl(testVal - reference, ulp_exp);

3
test_common/harness/imageHelpers.cpp
View File

@@ -690,9 +690,6 @@ int has_alpha(const cl_image_format *format)
_b ^= _a; \
_a ^= _b; \
} while (0)
#ifndef MAX
#define MAX(_a, _b) ((_a) > (_b) ? (_a) : (_b))
#endif
void get_max_sizes(
size_t *numberOfSizes, const int maxNumberOfSizes, size_t sizes[][3],

8
test_conformance/basic/test_enqueued_local_size.cpp
View File

@@ -14,13 +14,15 @@
// limitations under the License.
//
#include "harness/compat.h"
#include "harness/rounding_mode.h"
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <sys/types.h>
#include <sys/stat.h>
#include "harness/rounding_mode.h"
#include <algorithm>
#include "procs.h"
@@ -124,8 +126,8 @@ test_enqueued_local_size(cl_device_id device, cl_context context, cl_command_que
err = clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(max_wgs), &max_wgs, NULL);
test_error( err, "clGetDeviceInfo failed.");
localsize[0] = MIN(16, max_wgs);
localsize[1] = MIN(11, max_wgs / localsize[0]);
localsize[0] = std::min<size_t>(16, max_wgs);
localsize[1] = std::min<size_t>(11, max_wgs / localsize[0]);
// If we need to use uniform workgroups because non-uniform workgroups are
// not supported, round up to the next global size that is divisible by the
// local size.

22
test_conformance/buffers/test_sub_buffers.cpp
View File

@@ -15,6 +15,8 @@
//
#include "procs.h"
#include <algorithm>
// Design:
// To test sub buffers, we first create one main buffer. We then create several sub-buffers and
// queue Actions on each one. Each Action is encapsulated in a class so it can keep track of
@@ -101,13 +103,6 @@ public:
}
};
#ifndef MAX
#define MAX( _a, _b ) ( (_a) > (_b) ? (_a) : (_b) )
#endif
#ifndef MIN
#define MIN( _a, _b ) ( (_a) < (_b) ? (_a) : (_b) )
#endif
class CopyAction : public Action
{
public:
@@ -117,7 +112,8 @@ public:
virtual cl_int Execute( cl_context context, cl_command_queue queue, cl_char tag, SubBufferWrapper &buffer1, SubBufferWrapper &buffer2, cl_char *parentBufferState )
{
// Copy from sub-buffer 1 to sub-buffer 2
size_t size = get_random_size_t( 0, MIN( buffer1.mSize, buffer2.mSize ), GetRandSeed() );
size_t size = get_random_size_t(
0, std::min(buffer1.mSize, buffer2.mSize), GetRandSeed());
size_t startOffset = get_random_size_t( 0, buffer1.mSize - size, GetRandSeed() );
size_t endOffset = get_random_size_t( 0, buffer2.mSize - size, GetRandSeed() );
@@ -266,7 +262,11 @@ int test_sub_buffers_read_write_core( cl_context context, cl_command_queue queue
endRange = mainSize;
size_t offset = get_random_size_t( toStartFrom / addressAlign, endRange / addressAlign, Action::GetRandSeed() ) * addressAlign;
size_t size = get_random_size_t( 1, ( MIN( mainSize / 8, mainSize - offset ) ) / addressAlign, Action::GetRandSeed() ) * addressAlign;
size_t size =
get_random_size_t(
1, (std::min(mainSize / 8, mainSize - offset)) / addressAlign,
Action::GetRandSeed())
* addressAlign;
error = subBuffers[ numSubBuffers ].Allocate( mainBuffer, CL_MEM_READ_WRITE, offset, size );
test_error( error, "Unable to allocate sub buffer" );
@@ -443,7 +443,7 @@ int test_sub_buffers_read_write_dual_devices( cl_device_id deviceID, cl_context
error = get_reasonable_buffer_size( otherDevice, maxBuffer2 );
test_error( error, "Unable to get buffer size for secondary device" );
maxBuffer1 = MIN( maxBuffer1, maxBuffer2 );
maxBuffer1 = std::min(maxBuffer1, maxBuffer2);
cl_uint addressAlign1Bits, addressAlign2Bits;
error = clGetDeviceInfo( deviceID, CL_DEVICE_MEM_BASE_ADDR_ALIGN, sizeof( addressAlign1Bits ), &addressAlign1Bits, NULL );
@@ -452,7 +452,7 @@ int test_sub_buffers_read_write_dual_devices( cl_device_id deviceID, cl_context
error = clGetDeviceInfo( otherDevice, CL_DEVICE_MEM_BASE_ADDR_ALIGN, sizeof( addressAlign2Bits ), &addressAlign2Bits, NULL );
test_error( error, "Unable to get secondary device's address alignment" );
cl_uint addressAlign1 = MAX( addressAlign1Bits, addressAlign2Bits ) / 8;
cl_uint addressAlign1 = std::max(addressAlign1Bits, addressAlign2Bits) / 8;
// Finally time to run!
return test_sub_buffers_read_write_core( testingContext, queue1, queue2, maxBuffer1, addressAlign1 );

7
test_conformance/conversions/test_conversions.cpp
View File

@@ -47,6 +47,8 @@
#endif
#include <time.h>
#include <algorithm>
#include "Sleep.h"
#include "basic_test_conversions.h"
@@ -1003,7 +1005,8 @@ static int DoTest( cl_device_id device, Type outType, Type inType, SaturationMod
uint64_t i;
gTestCount++;
size_t blockCount = BUFFER_SIZE / MAX( gTypeSizes[ inType ], gTypeSizes[ outType ] );
size_t blockCount =
BUFFER_SIZE / std::max(gTypeSizes[inType], gTypeSizes[outType]);
size_t step = blockCount;
uint64_t lastCase = 1ULL << (8*gTypeSizes[ inType ]);
cl_event writeInputBuffer = NULL;
@@ -1078,7 +1081,7 @@ static int DoTest( cl_device_id device, Type outType, Type inType, SaturationMod
fflush(stdout);
}
cl_uint count = (uint32_t) MIN( blockCount, lastCase - i );
cl_uint count = (uint32_t)std::min((uint64_t)blockCount, lastCase - i);
writeInputBufferInfo.count = count;
// Crate a user event to represent the status of the reference value computation completion

3
test_conformance/device_execution/enqueue_ndrange.cpp
View File

@@ -18,6 +18,7 @@
#include "harness/testHarness.h"
#include "harness/typeWrappers.h"
#include <algorithm>
#include <vector>
#include "procs.h"
@@ -645,7 +646,7 @@ int test_enqueue_ndrange(cl_device_id device, cl_context context, cl_command_que
max_local_size = (max_local_size > MAX_GWS)? MAX_GWS: max_local_size;
if(gWimpyMode)
{
max_local_size = MIN(8, max_local_size);
max_local_size = std::min((size_t)8, max_local_size);
}
cl_uint num = 10;

3
test_conformance/device_execution/host_queue_order.cpp
View File

@@ -18,6 +18,7 @@
#include "harness/testHarness.h"
#include "harness/typeWrappers.h"
#include <algorithm>
#include <vector>
#include "procs.h"
@@ -124,7 +125,7 @@ int test_host_queue_order(cl_device_id device, cl_context context, cl_command_qu
cl_uint num = arr_size(result);
if( gWimpyMode )
{
num = MAX(num / 16, 4);
num = std::max(num / 16, 4U);
}
clMemWrapper res_mem;

7
test_conformance/half/Test_roundTrip.cpp
View File

@@ -14,6 +14,9 @@
// limitations under the License.
//
#include <string.h>
#include <algorithm>
#include "cl_utils.h"
#include "tests.h"
#include "harness/testHarness.h"
@@ -156,7 +159,7 @@ int test_roundTrip( cl_device_id device, cl_context context, cl_command_queue qu
}
// Figure out how many elements are in a work block
size_t elementSize = MAX( sizeof(cl_half), sizeof(cl_float));
size_t elementSize = std::max(sizeof(cl_half), sizeof(cl_float));
size_t blockCount = (size_t)getBufferSize(device) / elementSize; //elementSize is a power of two
uint64_t lastCase = 1ULL << (8*sizeof(cl_half)); // number of cl_half
size_t stride = blockCount;
@@ -168,7 +171,7 @@ int test_roundTrip( cl_device_id device, cl_context context, cl_command_queue qu
for( i = 0; i < (uint64_t)lastCase; i += stride )
{
count = (uint32_t) MIN( blockCount, lastCase - i );
count = (uint32_t)std::min((uint64_t)blockCount, lastCase - i);
//Init the input stream
uint16_t *p = (uint16_t *)gIn_half;

7
test_conformance/half/Test_vLoadHalf.cpp
View File

@@ -17,6 +17,9 @@
#include "harness/testHarness.h"
#include <string.h>
#include <algorithm>
#include "cl_utils.h"
#include "tests.h"
@@ -429,7 +432,7 @@ int Test_vLoadHalf_private( cl_device_id device, bool aligned )
}
// Figure out how many elements are in a work block
size_t elementSize = MAX( sizeof(cl_half), sizeof(cl_float));
size_t elementSize = std::max(sizeof(cl_half), sizeof(cl_float));
size_t blockCount = getBufferSize(device) / elementSize; // elementSize is power of 2
uint64_t lastCase = 1ULL << (8*sizeof(cl_half)); // number of things of size cl_half
@@ -447,7 +450,7 @@ int Test_vLoadHalf_private( cl_device_id device, bool aligned )
for( i = 0; i < (uint64_t)lastCase; i += blockCount )
{
count = (uint32_t) MIN( blockCount, lastCase - i );
count = (uint32_t)std::min((uint64_t)blockCount, lastCase - i);
//Init the input stream
uint16_t *p = (uint16_t *)gIn_half;

11
test_conformance/half/Test_vStoreHalf.cpp
View File

@@ -18,6 +18,9 @@
#include "harness/testHarness.h"
#include <string.h>
#include <algorithm>
#include "cl_utils.h"
#include "tests.h"
@@ -674,7 +677,7 @@ int Test_vStoreHalf_private( cl_device_id device, f2h referenceFunc, d2h doubleR
} // end for vector size
// Figure out how many elements are in a work block
size_t elementSize = MAX( sizeof(cl_ushort), sizeof(float));
size_t elementSize = std::max(sizeof(cl_ushort), sizeof(float));
size_t blockCount = BUFFER_SIZE / elementSize; // elementSize is power of 2
uint64_t lastCase = 1ULL << (8*sizeof(float)); // number of floats.
size_t stride = blockCount;
@@ -726,7 +729,7 @@ int Test_vStoreHalf_private( cl_device_id device, f2h referenceFunc, d2h doubleR
for( i = 0; i < lastCase; i += stride )
{
count = (cl_uint) MIN( blockCount, lastCase - i );
count = (cl_uint)std::min((uint64_t)blockCount, lastCase - i);
fref.i = i;
dref.i = i;
@@ -1272,7 +1275,7 @@ int Test_vStoreaHalf_private( cl_device_id device, f2h referenceFunc, d2h double
}
// Figure out how many elements are in a work block
size_t elementSize = MAX( sizeof(cl_ushort), sizeof(float));
size_t elementSize = std::max(sizeof(cl_ushort), sizeof(float));
size_t blockCount = BUFFER_SIZE / elementSize;
uint64_t lastCase = 1ULL << (8*sizeof(float));
size_t stride = blockCount;
@@ -1323,7 +1326,7 @@ int Test_vStoreaHalf_private( cl_device_id device, f2h referenceFunc, d2h double
for( i = 0; i < (uint64_t)lastCase; i += stride )
{
count = (cl_uint) MIN( blockCount, lastCase - i );
count = (cl_uint)std::min((uint64_t)blockCount, lastCase - i);
fref.i = i;
dref.i = i;

17
test_conformance/images/kernel_read_write/test_common.cpp
View File

@@ -16,6 +16,7 @@
#include "test_common.h"
#include <algorithm>
cl_sampler create_sampler(cl_context context, image_sampler_data *sdata, bool test_mipmaps, cl_int *error) {
cl_sampler sampler = nullptr;
@@ -934,13 +935,13 @@ int test_read_image(cl_context context, cl_command_queue queue,
{
err4 = 0.0f;
}
float maxErr1 = MAX(
float maxErr1 = std::max(
maxErr * maxPixel.p[0], FLT_MIN);
float maxErr2 = MAX(
float maxErr2 = std::max(
maxErr * maxPixel.p[1], FLT_MIN);
float maxErr3 = MAX(
float maxErr3 = std::max(
maxErr * maxPixel.p[2], FLT_MIN);
float maxErr4 = MAX(
float maxErr4 = std::max(
maxErr * maxPixel.p[3], FLT_MIN);
if (!(err1 <= maxErr1)
@@ -1039,16 +1040,16 @@ int test_read_image(cl_context context, cl_command_queue queue,
float err4 = ABS_ERROR(resultPtr[3],
expected[3]);
float maxErr1 =
MAX(maxErr * maxPixel.p[0],
std::max(maxErr * maxPixel.p[0],
FLT_MIN);
float maxErr2 =
MAX(maxErr * maxPixel.p[1],
std::max(maxErr * maxPixel.p[1],
FLT_MIN);
float maxErr3 =
MAX(maxErr * maxPixel.p[2],
std::max(maxErr * maxPixel.p[2],
FLT_MIN);
float maxErr4 =
MAX(maxErr * maxPixel.p[3],
std::max(maxErr * maxPixel.p[3],
FLT_MIN);

32
test_conformance/images/kernel_read_write/test_iterations.cpp
View File

@@ -16,6 +16,8 @@
#include "test_common.h"
#include <float.h>
#include <algorithm>
#if defined( __APPLE__ )
#include <signal.h>
#include <sys/signal.h>
@@ -434,7 +436,8 @@ int validate_image_2D_depth_results(void *imageValues, void *resultValues, doubl
float err1 = ABS_ERROR(resultPtr[0], expected[0]);
// Clamp to the minimum absolute error for the format
if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; }
float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN );
float maxErr1 =
std::max(maxErr * maxPixel.p[0], FLT_MIN);
// Check if the result matches.
if( ! (err1 <= maxErr1) )
@@ -484,7 +487,8 @@ int validate_image_2D_depth_results(void *imageValues, void *resultValues, doubl
imageSampler, expected, 0, &containsDenormals );
float err1 = ABS_ERROR(resultPtr[0], expected[0]);
float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN );
float maxErr1 =
std::max(maxErr * maxPixel.p[0], FLT_MIN);
if( ! (err1 <= maxErr1) )
@@ -598,10 +602,14 @@ int validate_image_2D_results(void *imageValues, void *resultValues, double form
if (err2 > 0 && err2 < formatAbsoluteError) { err2 = 0.0f; }
if (err3 > 0 && err3 < formatAbsoluteError) { err3 = 0.0f; }
if (err4 > 0 && err4 < formatAbsoluteError) { err4 = 0.0f; }
float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN );
float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN );
float maxErr3 = MAX( maxErr * maxPixel.p[2], FLT_MIN );
float maxErr4 = MAX( maxErr * maxPixel.p[3], FLT_MIN );
float maxErr1 =
std::max(maxErr * maxPixel.p[0], FLT_MIN);
float maxErr2 =
std::max(maxErr * maxPixel.p[1], FLT_MIN);
float maxErr3 =
std::max(maxErr * maxPixel.p[2], FLT_MIN);
float maxErr4 =
std::max(maxErr * maxPixel.p[3], FLT_MIN);
// Check if the result matches.
if( ! (err1 <= maxErr1) || ! (err2 <= maxErr2) ||
@@ -671,10 +679,14 @@ int validate_image_2D_results(void *imageValues, void *resultValues, double form
float err2 = ABS_ERROR(resultPtr[1], expected[1]);
float err3 = ABS_ERROR(resultPtr[2], expected[2]);
float err4 = ABS_ERROR(resultPtr[3], expected[3]);
float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN );
float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN );
float maxErr3 = MAX( maxErr * maxPixel.p[2], FLT_MIN );
float maxErr4 = MAX( maxErr * maxPixel.p[3], FLT_MIN );
float maxErr1 =
std::max(maxErr * maxPixel.p[0], FLT_MIN);
float maxErr2 =
std::max(maxErr * maxPixel.p[1], FLT_MIN);
float maxErr3 =
std::max(maxErr * maxPixel.p[2], FLT_MIN);
float maxErr4 =
std::max(maxErr * maxPixel.p[3], FLT_MIN);
if( ! (err1 <= maxErr1) || ! (err2 <= maxErr2) ||

26
test_conformance/images/kernel_read_write/test_read_1D.cpp
View File

@@ -17,6 +17,8 @@
#include "test_common.h"
#include <float.h>
#include <algorithm>
#if defined( __APPLE__ )
#include <signal.h>
#include <sys/signal.h>
@@ -669,10 +671,14 @@ int test_read_image_1D( cl_context context, cl_command_queue queue, cl_kernel ke
if (err2 > 0 && err2 < formatAbsoluteError) { err2 = 0.0f; }
if (err3 > 0 && err3 < formatAbsoluteError) { err3 = 0.0f; }
if (err4 > 0 && err4 < formatAbsoluteError) { err4 = 0.0f; }
float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN );
float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN );
float maxErr3 = MAX( maxErr * maxPixel.p[2], FLT_MIN );
float maxErr4 = MAX( maxErr * maxPixel.p[3], FLT_MIN );
float maxErr1 =
std::max(maxErr * maxPixel.p[0], FLT_MIN);
float maxErr2 =
std::max(maxErr * maxPixel.p[1], FLT_MIN);
float maxErr3 =
std::max(maxErr * maxPixel.p[2], FLT_MIN);
float maxErr4 =
std::max(maxErr * maxPixel.p[3], FLT_MIN);
// Check if the result matches.
if( ! (err1 <= maxErr1) || ! (err2 <= maxErr2) ||
@@ -732,10 +738,14 @@ int test_read_image_1D( cl_context context, cl_command_queue queue, cl_kernel ke
ABS_ERROR(resultPtr[2], expected[2]);
float err4 =
ABS_ERROR(resultPtr[3], expected[3]);
float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN );
float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN );
float maxErr3 = MAX( maxErr * maxPixel.p[2], FLT_MIN );
float maxErr4 = MAX( maxErr * maxPixel.p[3], FLT_MIN );
float maxErr1 =
std::max(maxErr * maxPixel.p[0], FLT_MIN);
float maxErr2 =
std::max(maxErr * maxPixel.p[1], FLT_MIN);
float maxErr3 =
std::max(maxErr * maxPixel.p[2], FLT_MIN);
float maxErr4 =
std::max(maxErr * maxPixel.p[3], FLT_MIN);
if( ! (err1 <= maxErr1) || ! (err2 <= maxErr2) ||

27
test_conformance/images/kernel_read_write/test_read_1D_array.cpp
View File

@@ -16,13 +16,14 @@
#include "test_common.h"
#include <float.h>
#include <algorithm>
#if defined( __APPLE__ )
#include <signal.h>
#include <sys/signal.h>
#include <setjmp.h>
#endif
const char *read1DArrayKernelSourcePattern =
"__kernel void sample_kernel( read_only image1d_array_t input,%s __global float *xOffsets, __global float *yOffsets, __global %s4 *results %s)\n"
"{\n"
@@ -772,10 +773,14 @@ int test_read_image_1D_array( cl_context context, cl_command_queue queue, cl_ker
if (err2 > 0 && err2 < formatAbsoluteError) { err2 = 0.0f; }
if (err3 > 0 && err3 < formatAbsoluteError) { err3 = 0.0f; }
if (err4 > 0 && err4 < formatAbsoluteError) { err4 = 0.0f; }
float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN );
float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN );
float maxErr3 = MAX( maxErr * maxPixel.p[2], FLT_MIN );
float maxErr4 = MAX( maxErr * maxPixel.p[3], FLT_MIN );
float maxErr1 =
std::max(maxErr * maxPixel.p[0], FLT_MIN);
float maxErr2 =
std::max(maxErr * maxPixel.p[1], FLT_MIN);
float maxErr3 =
std::max(maxErr * maxPixel.p[2], FLT_MIN);
float maxErr4 =
std::max(maxErr * maxPixel.p[3], FLT_MIN);
// Check if the result matches.
if( ! (err1 <= maxErr1) || ! (err2 <= maxErr2) ||
@@ -838,10 +843,14 @@ int test_read_image_1D_array( cl_context context, cl_command_queue queue, cl_ker
ABS_ERROR(resultPtr[2], expected[2]);
float err4 =
ABS_ERROR(resultPtr[3], expected[3]);
float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN );
float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN );
float maxErr3 = MAX( maxErr * maxPixel.p[2], FLT_MIN );
float maxErr4 = MAX( maxErr * maxPixel.p[3], FLT_MIN );
float maxErr1 =
std::max(maxErr * maxPixel.p[0], FLT_MIN);
float maxErr2 =
std::max(maxErr * maxPixel.p[1], FLT_MIN);
float maxErr3 =
std::max(maxErr * maxPixel.p[2], FLT_MIN);
float maxErr4 =
std::max(maxErr * maxPixel.p[3], FLT_MIN);
if( ! (err1 <= maxErr1) || ! (err2 <= maxErr2) ||

32
test_conformance/images/kernel_read_write/test_read_2D_array.cpp
View File

@@ -16,6 +16,8 @@
#include "test_common.h"
#include <float.h>
#include <algorithm>
// Utility function to clamp down image sizes for certain tests to avoid
// using too much memory.
static size_t reduceImageSizeRange(size_t maxDimSize) {
@@ -617,7 +619,8 @@ int test_read_image_2D_array( cl_context context, cl_command_queue queue, cl_ker
ABS_ERROR(resultPtr[0], expected[0]);
// Clamp to the minimum absolute error for the format
if (err1 > 0 && err1 < formatAbsoluteError) { err1 = 0.0f; }
float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN );
float maxErr1 = std::max(
maxErr * maxPixel.p[0], FLT_MIN);
if( ! (err1 <= maxErr1) )
{
@@ -661,7 +664,8 @@ int test_read_image_2D_array( cl_context context, cl_command_queue queue, cl_ker
float err1 = ABS_ERROR(resultPtr[0],
expected[0]);
float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN );
float maxErr1 = std::max(
maxErr * maxPixel.p[0], FLT_MIN);
if( ! (err1 <= maxErr1) )
@@ -942,10 +946,14 @@ int test_read_image_2D_array( cl_context context, cl_command_queue queue, cl_ker
if (err2 > 0 && err2 < formatAbsoluteError) { err2 = 0.0f; }
if (err3 > 0 && err3 < formatAbsoluteError) { err3 = 0.0f; }
if (err4 > 0 && err4 < formatAbsoluteError) { err4 = 0.0f; }
float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN );
float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN );
float maxErr3 = MAX( maxErr * maxPixel.p[2], FLT_MIN );
float maxErr4 = MAX( maxErr * maxPixel.p[3], FLT_MIN );
float maxErr1 = std::max(
maxErr * maxPixel.p[0], FLT_MIN);
float maxErr2 = std::max(
maxErr * maxPixel.p[1], FLT_MIN);
float maxErr3 = std::max(
maxErr * maxPixel.p[2], FLT_MIN);
float maxErr4 = std::max(
maxErr * maxPixel.p[3], FLT_MIN);
if( ! (err1 <= maxErr1) || ! (err2 <= maxErr2) || ! (err3 <= maxErr3) || ! (err4 <= maxErr4) )
{
@@ -1004,10 +1012,14 @@ int test_read_image_2D_array( cl_context context, cl_command_queue queue, cl_ker
expected[2]);
float err4 = ABS_ERROR(resultPtr[3],
expected[3]);
float maxErr1 = MAX( maxErr * maxPixel.p[0], FLT_MIN );
float maxErr2 = MAX( maxErr * maxPixel.p[1], FLT_MIN );
float maxErr3 = MAX( maxErr * maxPixel.p[2], FLT_MIN );
float maxErr4 = MAX( maxErr * maxPixel.p[3], FLT_MIN );
float maxErr1 = std::max(
maxErr * maxPixel.p[0], FLT_MIN);
float maxErr2 = std::max(
maxErr * maxPixel.p[1], FLT_MIN);
float maxErr3 = std::max(
maxErr * maxPixel.p[2], FLT_MIN);
float maxErr4 = std::max(
maxErr * maxPixel.p[3], FLT_MIN);
if( ! (err1 <= maxErr1) || ! (err2 <= maxErr2) || ! (err3 <= maxErr3) || ! (err4 <= maxErr4) )

31
test_conformance/integer_ops/test_add_sat.cpp
View File

@@ -21,27 +21,18 @@
#include <sys/types.h>
#include <sys/stat.h>
#include <algorithm>
#include "procs.h"
#define UCHAR_MIN 0
#define USHRT_MIN 0
#define UINT_MIN 0
#ifndef MAX
#define MAX( _a, _b ) ( (_a) > (_b) ? (_a) : (_b) )
#endif
#ifndef MIN
#define MIN( _a, _b ) ( (_a) < (_b) ? (_a) : (_b) )
#endif
static int verify_addsat_char( const cl_char *inA, const cl_char *inB, const cl_char *outptr, int n, const char *sizeName, int vecSize )
{
int i;
for( i = 0; i < n; i++ )
{
cl_int r = (cl_int) inA[i] + (cl_int) inB[i];
r = MAX( r, CL_CHAR_MIN );
r = MIN( r, CL_CHAR_MAX );
r = std::max(r, CL_CHAR_MIN);
r = std::min(r, CL_CHAR_MAX);
if( r != outptr[i] )
{ log_info( "\n%d) Failure for add_sat( (char%s) 0x%2.2x, (char%s) 0x%2.2x) = *0x%2.2x vs 0x%2.2x\n", i, sizeName, inA[i], sizeName, inB[i], r, outptr[i] ); return -1; }
@@ -55,9 +46,9 @@ static int verify_addsat_uchar( const cl_uchar *inA, const cl_uchar *inB, const
for( i = 0; i < n; i++ )
{
cl_int r = (int) inA[i] + (int) inB[i];
r = MAX( r, 0 );
r = MIN( r, CL_UCHAR_MAX );
if( r != outptr[i] )
r = std::max(r, 0);
r = std::min(r, CL_UCHAR_MAX);
if (r != outptr[i])
{ log_info( "\n%d) Failure for add_sat( (uchar%s) 0x%2.2x, (uchar%s) 0x%2.2x) = *0x%2.2x vs 0x%2.2x\n", i, sizeName, inA[i], sizeName, inB[i], r, outptr[i] ); return -1; }
}
return 0;
@@ -69,8 +60,8 @@ static int verify_addsat_short( const cl_short *inA, const cl_short *inB, const
for( i = 0; i < n; i++ )
{
cl_int r = (cl_int) inA[i] + (cl_int) inB[i];
r = MAX( r, CL_SHRT_MIN );
r = MIN( r, CL_SHRT_MAX );
r = std::max(r, CL_SHRT_MIN);
r = std::min(r, CL_SHRT_MAX);
if( r != outptr[i] )
{ log_info( "\n%d) Failure for add_sat( (short%s) 0x%4.4x, (short%s) 0x%4.4x) = *0x%4.4x vs 0x%4.4x\n", i, sizeName, inA[i], sizeName, inB[i], r, outptr[i] ); return -1; }
@@ -84,8 +75,8 @@ static int verify_addsat_ushort( const cl_ushort *inA, const cl_ushort *inB, con
for( i = 0; i < n; i++ )
{
cl_int r = (cl_int) inA[i] + (cl_int) inB[i];
r = MAX( r, 0 );
r = MIN( r, CL_USHRT_MAX );
r = std::max(r, 0);
r = std::min(r, CL_USHRT_MAX);
if( r != outptr[i] )
{ log_info( "\n%d) Failure for add_sat( (ushort%s) 0x%4.4x, (ushort%s) 0x%4.4x) = *0x%4.4x vs 0x%4.4x\n", i, sizeName, inA[i], sizeName, inB[i], r, outptr[i] ); return -1; }

60
test_conformance/integer_ops/test_integers.cpp
View File

@@ -16,14 +16,9 @@
#include "testBase.h"
#include "harness/conversions.h"
#define TEST_SIZE 512
#include <algorithm>
#ifndef MIN
#define MIN( _a, _b ) ((_a) < (_b) ? (_a) : (_b))
#endif
#ifndef MAX
#define MAX( _a, _b ) ((_a) > (_b) ? (_a) : (_b))
#endif
#define TEST_SIZE 512
const char *singleParamIntegerKernelSourcePattern =
"__kernel void sample_test(__global %s *sourceA, __global %s *destValues)\n"
@@ -1512,19 +1507,20 @@ bool verify_integer_clamp( void *sourceA, void *sourceB, void *sourceC, void *de
switch( vecAType )
{
case kULong:
((cl_ulong*) destination)[0] = MAX(MIN(valueA, valueC), valueB);
((cl_ulong *)destination)[0] =
std::max(std::min(valueA, valueC), valueB);
break;
case kUInt:
((cl_uint*) destination)[0] = (cl_uint)
(MAX(MIN(valueA, valueC), valueB));
((cl_uint *)destination)[0] =
(cl_uint)(std::max(std::min(valueA, valueC), valueB));
break;
case kUShort:
((cl_ushort*) destination)[0] = (cl_ushort)
(MAX(MIN(valueA, valueC), valueB));
((cl_ushort *)destination)[0] =
(cl_ushort)(std::max(std::min(valueA, valueC), valueB));
break;
case kUChar:
((cl_uchar*) destination)[0] = (cl_uchar)
(MAX(MIN(valueA, valueC), valueB));
((cl_uchar *)destination)[0] =
(cl_uchar)(std::max(std::min(valueA, valueC), valueB));
break;
default:
//error -- should never get here
@@ -1576,19 +1572,20 @@ bool verify_integer_clamp( void *sourceA, void *sourceB, void *sourceC, void *de
switch( vecAType )
{
case kLong:
((cl_long*) destination)[0] = MAX(MIN(valueA, valueC), valueB);
((cl_long *)destination)[0] =
std::max(std::min(valueA, valueC), valueB);
break;
case kInt:
((cl_int*) destination)[0] = (cl_int)
(MAX(MIN(valueA, valueC), valueB));
((cl_int *)destination)[0] =
(cl_int)(std::max(std::min(valueA, valueC), valueB));
break;
case kShort:
((cl_short*) destination)[0] = (cl_short)
(MAX(MIN(valueA, valueC), valueB));
((cl_short *)destination)[0] =
(cl_short)(std::max(std::min(valueA, valueC), valueB));
break;
case kChar:
((cl_char*) destination)[0] = (cl_char)
(MAX(MIN(valueA, valueC), valueB));
((cl_char *)destination)[0] =
(cl_char)(std::max(std::min(valueA, valueC), valueB));
break;
default:
//error -- should never get here
@@ -1654,13 +1651,16 @@ bool verify_integer_mad_sat( void *sourceA, void *sourceB, void *sourceC, void *
((cl_ulong*) destination)[0] = multLo;
break;
case kUInt:
((cl_uint*) destination)[0] = (cl_uint) MIN( multLo, (cl_ulong) CL_UINT_MAX );
((cl_uint *)destination)[0] =
(cl_uint)std::min(multLo, (cl_ulong)CL_UINT_MAX);
break;
case kUShort:
((cl_ushort*) destination)[0] = (cl_ushort) MIN( multLo, (cl_ulong) CL_USHRT_MAX );
((cl_ushort *)destination)[0] =
(cl_ushort)std::min(multLo, (cl_ulong)CL_USHRT_MAX);
break;
case kUChar:
((cl_uchar*) destination)[0] = (cl_uchar) MIN( multLo, (cl_ulong) CL_UCHAR_MAX );
((cl_uchar *)destination)[0] =
(cl_uchar)std::min(multLo, (cl_ulong)CL_UCHAR_MAX);
break;
default:
//error -- should never get here
@@ -1744,18 +1744,18 @@ bool verify_integer_mad_sat( void *sourceA, void *sourceB, void *sourceC, void *
((cl_long*) destination)[0] = result;
break;
case kInt:
result = MIN( result, (cl_long) CL_INT_MAX );
result = MAX( result, (cl_long) CL_INT_MIN );
result = std::min(result, (cl_long)CL_INT_MAX);
result = std::max(result, (cl_long)CL_INT_MIN);
((cl_int*) destination)[0] = (cl_int) result;
break;
case kShort:
result = MIN( result, (cl_long) CL_SHRT_MAX );
result = MAX( result, (cl_long) CL_SHRT_MIN );
result = std::min(result, (cl_long)CL_SHRT_MAX);
result = std::max(result, (cl_long)CL_SHRT_MIN);
((cl_short*) destination)[0] = (cl_short) result;
break;
case kChar:
result = MIN( result, (cl_long) CL_CHAR_MAX );
result = MAX( result, (cl_long) CL_CHAR_MIN );
result = std::min(result, (cl_long)CL_CHAR_MAX);
result = std::max(result, (cl_long)CL_CHAR_MIN);
((cl_char*) destination)[0] = (cl_char) result;
break;
default:

32
test_conformance/integer_ops/test_sub_sat.cpp
View File

@@ -21,28 +21,18 @@
#include <sys/types.h>
#include <sys/stat.h>
#include <algorithm>
#include "procs.h"
#define UCHAR_MIN 0
#define USHRT_MIN 0
#define UINT_MIN 0
#ifndef MAX
#define MAX( _a, _b ) ( (_a) > (_b) ? (_a) : (_b) )
#endif
#ifndef MIN
#define MIN( _a, _b ) ( (_a) < (_b) ? (_a) : (_b) )
#endif
static int verify_subsat_char( const cl_char *inA, const cl_char *inB, const cl_char *outptr, int n, const char *sizeName, int vecSize )
{
int i;
for( i = 0; i < n; i++ )
{
cl_int r = (cl_int) inA[i] - (cl_int) inB[i];
r = MAX( r, CL_CHAR_MIN );
r = MIN( r, CL_CHAR_MAX );
r = std::max(r, CL_CHAR_MIN);
r = std::min(r, CL_CHAR_MAX);
if( r != outptr[i] )
{ log_info( "\n%d) Failure for sub_sat( (char%s) 0x%2.2x, (char%s) 0x%2.2x) = *0x%2.2x vs 0x%2.2x\n", i, sizeName, inA[i], sizeName, inB[i], r, outptr[i] ); return -1; }
@@ -56,9 +46,9 @@ static int verify_subsat_uchar( const cl_uchar *inA, const cl_uchar *inB, const
for( i = 0; i < n; i++ )
{
cl_int r = (cl_int) inA[i] - (cl_int) inB[i];
r = MAX( r, 0 );
r = MIN( r, CL_UCHAR_MAX );
if( r != outptr[i] )
r = std::max(r, 0);
r = std::min(r, CL_UCHAR_MAX);
if (r != outptr[i])
{ log_info( "\n%d) Failure for sub_sat( (uchar%s) 0x%2.2x, (uchar%s) 0x%2.2x) = *0x%2.2x vs 0x%2.2x\n", i, sizeName, inA[i], sizeName, inB[i], r, outptr[i] ); return -1; }
}
return 0;
@@ -70,8 +60,8 @@ static int verify_subsat_short( const cl_short *inA, const cl_short *inB, const
for( i = 0; i < n; i++ )
{
cl_int r = (cl_int) inA[i] - (cl_int) inB[i];
r = MAX( r, CL_SHRT_MIN );
r = MIN( r, CL_SHRT_MAX );
r = std::max(r, CL_SHRT_MIN);
r = std::min(r, CL_SHRT_MAX);
if( r != outptr[i] )
{ log_info( "\n%d) Failure for sub_sat( (short%s) 0x%4.4x, (short%s) 0x%4.4x) = *0x%4.4x vs 0x%4.4x\n", i, sizeName, inA[i], sizeName, inB[i], r, outptr[i] ); return -1; }
@@ -85,8 +75,8 @@ static int verify_subsat_ushort( const cl_ushort *inA, const cl_ushort *inB, con
for( i = 0; i < n; i++ )
{
cl_int r = (cl_int) inA[i] - (cl_int) inB[i];
r = MAX( r, 0 );
r = MIN( r, CL_USHRT_MAX );
r = std::max(r, 0);
r = std::min(r, CL_USHRT_MAX);
if( r != outptr[i] )
{ log_info( "\n%d) Failure for sub_sat( (ushort%s) 0x%4.4x, (ushort%s) 0x%4.4x) = *0x%4.4x vs 0x%4.4x\n", i, sizeName, inA[i], sizeName, inB[i], r, outptr[i] ); return -1; }

2
test_conformance/integer_ops/test_unary_ops.cpp
View File

@@ -107,7 +107,7 @@ int test_unary_op( cl_command_queue queue, cl_context context, OpKonstants which
// For sub ops, the min control value is 2. Otherwise, it's 0
controlData[ i ] |= 0x02;
else if( whichOp == kIncrement )
// For addition ops, the MAX control value is 1. Otherwise, it's 3
// For addition ops, the max control value is 1. Otherwise, it's 3
controlData[ i ] &= ~0x02;
}
streams[1] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,

3
test_conformance/math_brute_force/macro_binary_double.cpp
View File

@@ -496,7 +496,8 @@ cl_int Test(cl_uint job_id, cl_uint thread_id, void *data)
}
for (auto k = MAX(1, gMinVectorSizeIndex); k < gMaxVectorSizeIndex; k++)
for (auto k = std::max(1U, gMinVectorSizeIndex);
k < gMaxVectorSizeIndex; k++)
{
q = (cl_long *)out[k];
// If we aren't getting the correctly rounded result

3
test_conformance/math_brute_force/macro_binary_float.cpp
View File

@@ -485,7 +485,8 @@ cl_int Test(cl_uint job_id, cl_uint thread_id, void *data)
goto exit;
}
for (auto k = MAX(1, gMinVectorSizeIndex); k < gMaxVectorSizeIndex; k++)
for (auto k = std::max(1U, gMinVectorSizeIndex);
k < gMaxVectorSizeIndex; k++)
{
q = out[k];
// If we aren't getting the correctly rounded result

3
test_conformance/math_brute_force/macro_unary_double.cpp
View File

@@ -304,7 +304,8 @@ cl_int Test(cl_uint job_id, cl_uint thread_id, void *data)
}
for (auto k = MAX(1, gMinVectorSizeIndex); k < gMaxVectorSizeIndex; k++)
for (auto k = std::max(1U, gMinVectorSizeIndex);
k < gMaxVectorSizeIndex; k++)
{
q = out[k];
// If we aren't getting the correctly rounded result

4
test_conformance/math_brute_force/macro_unary_float.cpp
View File

@@ -309,8 +309,8 @@ cl_int Test(cl_uint job_id, cl_uint thread_id, void *data)
}
for (auto k = MAX(1, gMinVectorSizeIndex); k < gMaxVectorSizeIndex;
k++)
for (auto k = std::max(1U, gMinVectorSizeIndex);
k < gMaxVectorSizeIndex; k++)
{
q = out[k];
// If we aren't getting the correctly rounded result

5
test_conformance/math_brute_force/main.cpp
View File

@@ -18,6 +18,7 @@
#include "sleep.h"
#include "utility.h"
#include <algorithm>
#include <cstdio>
#include <cstdlib>
#include <ctime>
@@ -1239,7 +1240,7 @@ float Bruteforce_Ulp_Error_Double(double test, long double reference)
// The unbiased exponent of the ulp unit place
int ulp_exp =
DBL_MANT_DIG - 1 - MAX(ilogbl(reference), DBL_MIN_EXP - 1);
DBL_MANT_DIG - 1 - std::max(ilogbl(reference), DBL_MIN_EXP - 1);
// Scale the exponent of the error
float result = (float)scalbnl(testVal - reference, ulp_exp);
@@ -1255,7 +1256,7 @@ float Bruteforce_Ulp_Error_Double(double test, long double reference)
// reference is a normal power of two or a zero
// The unbiased exponent of the ulp unit place
int ulp_exp =
DBL_MANT_DIG - 1 - MAX(ilogbl(reference) - 1, DBL_MIN_EXP - 1);
DBL_MANT_DIG - 1 - std::max(ilogbl(reference) - 1, DBL_MIN_EXP - 1);
// allow correctly rounded results to pass through unmolested. (We might add
// error to it below.) There is something of a performance optimization here

12
test_conformance/profiling/execute.cpp
View File

@@ -21,6 +21,8 @@
#include <sys/types.h>
#include <sys/stat.h>
#include <algorithm>
#include "procs.h"
#include "harness/testHarness.h"
#include "harness/errorHelpers.h"
@@ -29,12 +31,6 @@
typedef unsigned char uchar;
#endif
#undef MIN
#define MIN(x,y) ( (x) < (y) ? (x) : (y) )
#undef MAX
#define MAX(x,y) ( (x) > (y) ? (x) : (y) )
//#define CREATE_OUTPUT 1
extern int writePPM( const char *filename, uchar *buf, int xsize, int ysize );
@@ -73,8 +69,8 @@ static const char *image_filter_src =
static void read_imagef( int x, int y, int w, int h, int nChannels, uchar *src, float *srcRgb )
{
// clamp the coords
int x0 = MIN( MAX( x, 0 ), w - 1 );
int y0 = MIN( MAX( y, 0 ), h - 1 );
int x0 = std::min(std::max(x, 0), w - 1);
int y0 = std::min(std::max(y, 0), h - 1);
// get tine index
int indx = ( y0 * w + x0 ) * nChannels;

6
test_conformance/workgroups/test_wg_broadcast.cpp
View File

@@ -20,6 +20,8 @@
#include <sys/types.h>
#include <sys/stat.h>
#include <algorithm>
#include "procs.h"
@@ -310,7 +312,7 @@ test_work_group_broadcast_2D(cl_device_id device, cl_context context, cl_command
localsize[0] = localsize[1] = 1;
}
num_workgroups = MAX(n_elems/wg_size[0], 16);
num_workgroups = std::max(n_elems / wg_size[0], (size_t)16);
globalsize[0] = num_workgroups * localsize[0];
globalsize[1] = num_workgroups * localsize[1];
num_elements = globalsize[0] * globalsize[1];
@@ -437,7 +439,7 @@ test_work_group_broadcast_3D(cl_device_id device, cl_context context, cl_command
localsize[0] = localsize[1] = localsize[2] = 1;
}
num_workgroups = MAX(n_elems/wg_size[0], 8);
num_workgroups = std::max(n_elems / wg_size[0], (size_t)8);
globalsize[0] = num_workgroups * localsize[0];
globalsize[1] = num_workgroups * localsize[1];
globalsize[2] = num_workgroups * localsize[2];

11
test_conformance/workgroups/test_wg_scan_exclusive_max.cpp
View File

@@ -20,8 +20,9 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include <algorithm>
#include "procs.h"
const char *wg_scan_exclusive_max_kernel_code_int =
"__kernel void test_wg_scan_exclusive_max_int(global int *input, global int *output)\n"
@@ -79,7 +80,7 @@ verify_wg_scan_exclusive_max_int(int *inptr, int *outptr, size_t n, size_t wg_si
log_info("work_group_scan_exclusive_max int: Error at %u: expected = %d, got = %d\n", (unsigned int)(j+i), max_, outptr[j+i]);
return -1;
}
max_ = MAX(inptr[j+i], max_);
max_ = std::max(inptr[j + i], max_);
}
}
@@ -103,7 +104,7 @@ verify_wg_scan_exclusive_max_uint(unsigned int *inptr, unsigned int *outptr, siz
log_info("work_group_scan_exclusive_max int: Error at %u: expected = %u, got = %u\n", (unsigned int)(j+i), max_, outptr[j+i]);
return -1;
}
max_ = MAX(inptr[j+i], max_);
max_ = std::max(inptr[j + i], max_);
}
}
@@ -127,7 +128,7 @@ verify_wg_scan_exclusive_max_long(cl_long *inptr, cl_long *outptr, size_t n, siz
log_info("work_group_scan_exclusive_max long: Error at %u: expected = %lld, got = %lld\n", (unsigned int)(j+i), max_, outptr[j+i]);
return -1;
}
max_ = MAX(inptr[j+i], max_);
max_ = std::max(inptr[j + i], max_);
}
}
@@ -151,7 +152,7 @@ verify_wg_scan_exclusive_max_ulong(cl_ulong *inptr, cl_ulong *outptr, size_t n,
log_info("work_group_scan_exclusive_max ulong: Error at %u: expected = %llu, got = %llu\n", (unsigned int)(j+i), max_, outptr[j+i]);
return -1;
}
max_ = MAX(inptr[j+i], max_);
max_ = std::max(inptr[j + i], max_);
}
}

11
test_conformance/workgroups/test_wg_scan_exclusive_min.cpp
View File

@@ -20,8 +20,9 @@
#include <sys/types.h>
#include <sys/stat.h>
#include "procs.h"
#include <algorithm>
#include "procs.h"
const char *wg_scan_exclusive_min_kernel_code_int =
"__kernel void test_wg_scan_exclusive_min_int(global int *input, global int *output)\n"
@@ -80,7 +81,7 @@ verify_wg_scan_exclusive_min_int(int *inptr, int *outptr, size_t n, size_t wg_si
log_info("work_group_scan_exclusive_min int: Error at %u: expected = %d, got = %d\n", (unsigned int)(j+i), min_, outptr[j+i]);
return -1;
}
min_ = MIN(inptr[j+i], min_);
min_ = std::min(inptr[j + i], min_);
}
}
@@ -104,7 +105,7 @@ verify_wg_scan_exclusive_min_uint(unsigned int *inptr, unsigned int *outptr, siz
log_info("work_group_scan_exclusive_min int: Error at %u: expected = %u, got = %u\n", j+i, min_, outptr[j+i]);
return -1;
}
min_ = MIN(inptr[j+i], min_);
min_ = std::min(inptr[j + i], min_);
}
}
@@ -128,7 +129,7 @@ verify_wg_scan_exclusive_min_long(cl_long *inptr, cl_long *outptr, size_t n, siz
log_info("work_group_scan_exclusive_min long: Error at %u: expected = %lld, got = %lld\n", (unsigned int)(j+i), min_, outptr[j+i]);
return -1;
}
min_ = MIN(inptr[j+i], min_);
min_ = std::min(inptr[j + i], min_);
}
}
@@ -152,7 +153,7 @@ verify_wg_scan_exclusive_min_ulong(cl_ulong *inptr, cl_ulong *outptr, size_t n,
log_info("work_group_scan_exclusive_min ulong: Error at %u: expected = %llu, got = %llu\n", (unsigned int)(j+i), min_, outptr[j+i]);
return -1;
}
min_ = MIN(inptr[j+i], min_);
min_ = std::min(inptr[j + i], min_);
}
}

10
test_conformance/workgroups/test_wg_scan_inclusive_max.cpp
View File

@@ -20,6 +20,8 @@
#include <sys/types.h>
#include <sys/stat.h>
#include <algorithm>
#include "procs.h"
@@ -75,7 +77,7 @@ verify_wg_scan_inclusive_max_int(int *inptr, int *outptr, size_t n, size_t wg_si
m = wg_size;
for (i = 0; i < m; ++i) {
max_ = MAX(inptr[j+i], max_);
max_ = std::max(inptr[j + i], max_);
if (outptr[j+i] != max_) {
log_info("work_group_scan_inclusive_max int: Error at %u: expected = %d, got = %d\n", (unsigned int)(j+i), max_, outptr[j+i]);
return -1;
@@ -99,7 +101,7 @@ verify_wg_scan_inclusive_max_uint(unsigned int *inptr, unsigned int *outptr, siz
m = wg_size;
for (i = 0; i < m; ++i) {
max_ = MAX(inptr[j+i], max_);
max_ = std::max(inptr[j + i], max_);
if (outptr[j+i] != max_) {
log_info("work_group_scan_inclusive_max int: Error at %lu: expected = %u, got = %u\n", (unsigned long)(j+i), max_, outptr[j+i]);
return -1;
@@ -123,7 +125,7 @@ verify_wg_scan_inclusive_max_long(cl_long *inptr, cl_long *outptr, size_t n, siz
m = wg_size;
for (i = 0; i < m; ++i) {
max_ = MAX(inptr[j+i], max_);
max_ = std::max(inptr[j + i], max_);
if (outptr[j+i] != max_) {
log_info("work_group_scan_inclusive_max long: Error at %u: expected = %lld, got = %lld\n", (unsigned int)(j+i), max_, outptr[j+i]);
return -1;
@@ -147,7 +149,7 @@ verify_wg_scan_inclusive_max_ulong(cl_ulong *inptr, cl_ulong *outptr, size_t n,
m = wg_size;
for (i = 0; i < m; ++i) {
max_ = MAX(inptr[j+i], max_);
max_ = std::max(inptr[j + i], max_);
if (outptr[j+i] != max_) {
log_info("work_group_scan_inclusive_max ulong: Error at %u: expected = %llu, got = %llu\n", (unsigned int)(j+i), max_, outptr[j+i]);
return -1;

10
test_conformance/workgroups/test_wg_scan_inclusive_min.cpp
View File

@@ -20,6 +20,8 @@
#include <sys/types.h>
#include <sys/stat.h>
#include <algorithm>
#include "procs.h"
@@ -75,7 +77,7 @@ verify_wg_scan_inclusive_min_int(int *inptr, int *outptr, size_t n, size_t wg_si
m = wg_size;
for (i = 0; i < m; ++i) {
min_ = MIN(inptr[j+i], min_);
min_ = std::min(inptr[j + i], min_);
if (outptr[j+i] != min_) {
log_info("work_group_scan_inclusive_min int: Error at %u: expected = %d, got = %d\n", (unsigned int)(j+i), min_, outptr[j+i]);
return -1;
@@ -99,7 +101,7 @@ verify_wg_scan_inclusive_min_uint(unsigned int *inptr, unsigned int *outptr, siz
m = wg_size;
for (i = 0; i < m; ++i) {
min_ = MIN(inptr[j+i], min_);
min_ = std::min(inptr[j + i], min_);
if (outptr[j+i] != min_) {
log_info("work_group_scan_inclusive_min int: Error at %u: expected = %u, got = %u\n", (unsigned int)(j+i), min_, outptr[j+i]);
return -1;
@@ -123,7 +125,7 @@ verify_wg_scan_inclusive_min_long(cl_long *inptr, cl_long *outptr, size_t n, siz
m = wg_size;
for (i = 0; i < m; ++i) {
min_ = MIN(inptr[j+i], min_);
min_ = std::min(inptr[j + i], min_);
if (outptr[j+i] != min_) {
log_info("work_group_scan_inclusive_min long: Error at %u: expected = %lld, got = %lld\n", (unsigned int)(j+i), min_, outptr[j+i]);
return -1;
@@ -147,7 +149,7 @@ verify_wg_scan_inclusive_min_ulong(cl_ulong *inptr, cl_ulong *outptr, size_t n,
m = wg_size;
for (i = 0; i < m; ++i) {
min_ = MIN(inptr[j+i], min_);
min_ = std::min(inptr[j + i], min_);
if (outptr[j+i] != min_) {
log_info("work_group_scan_inclusive_min ulong: Error at %u: expected = %llu, got = %llu\n", (unsigned int)(j+i), min_, outptr[j+i]);
return -1;