Skip to content

Latest commit

 

History

History
654 lines (545 loc) · 15.6 KB

individual-functions.asciidoc

File metadata and controls

654 lines (545 loc) · 15.6 KB

Test plan for individual functions (2020 update)

This is a test plan for the changes in APIs described in SYCL 2020 compared to SYCL 1.2.1 relating to mathematical, integer, common, geometric, and relational functions.

Estimated development time is six days.

1. Testing scope

Tests intend to extend partial to full coverage for:

  • 4.17.2. Function objects

  • 4.17.5. Math functions

  • 4.17.6. Integer functions

  • 4.17.7. Common functions

  • 4.17.8. Geometric functions

  • 4.17.9. Relational functions

1.1. Device coverage

Device coverage does not change. All of the tests described below are performed only on the default device that is selected on the CTS command line.

1.2. Types coverage

Most of the tests are extensions of the existing tests to new marray types. The system of types coverage is coded in Python code generators therefore these type generators will be extended to cover full Table 165. Generic type name description, which serves as a description for all valid types of parameters to kernel functions.

More specifically, new marray types will be introduced with 1, 2, 3, 4, 5, and 17 elements and they will be added to the following generic type names (this will extend also more generic names):

Generic type name Description
floatn

float{n}, mfloat{n}, marray<{N},float>

doublen

double{n}, mdouble{n}, marray<{N},double>

halfn

half{n}, mhalf{n}, marray<{N},half>

mgenfloat

marray<float,{N}>, marray<double,{N}>, marray<half,{N}>

gengeofloat

float, float2, float3, float4, mfloat2, mfloat3, mfloat4

gengeodouble

double, double2, double3, double4, mdouble2, mdouble3, mdouble4

charn

char{n}, mchar{n}, marray<{N},char>

scharn

schar{n}, mschar{n}, marray<{N},signed char>

ucharn

uchar{n}, muchar{n}, marray<{N},unsigned char>

shortn

short{n}, mshort{n}, marray<{N},short>

ushortn

ushort{n}, mushort{n}, marray<{N},unsigned short>

uintn

uint{n}, muint{n}, marray<{N},unsigned int>

intn

int{n}, mint{n}, marray<{N},int>

ulongn

ulong{n}, mulong{n}, marray<{N},unsigned long int>

longn

long{n}, mlong{n}, marray<{N},long int>

ulonglongn

ulonglong{n}, mulonglong{n}, marray<{N},unsigned long long int>

longlongn

longlong{n}, mlonglong{n}, marray<{N},long long int>

migeninteger

marray<signed char,{N}>, marray<short,{N}>, marray<int,{N}>, marray<long,{N}>, marray<long long,{N}>

mugeninteger

marray<unsigned char,{N}>, marray<unsigned short,{N}>, marray<unsigned int,{N}>, marray<unsigned long int,{N}>, marray<unsigned long long int,{N}>

mgeninteger

marray<char,{N}>, migeninteger, mugeninteger

mgentype

mgenfloat, mgeninteger

booln

marray<{N},bool>

genbool

bool, booln

If the device has aspect::fp64, double type is tested, and if the device has aspect::fp16, sycl::half type is tested.

Reference templated implementations for marray functions will be added to sycl-cts/util/math_reference.h.

2. Tests

2.1. 4.17.2. Function objects

The function objects are already tested in CTS in reduction tests. The only aspect non-tested is their <void> specializations that are transparent function objects which deduce their parameter types and return type. Then the test will check that C++ conversion and promotion rules are satisfied for <void> specializations of the function objects.

2.2. 4.17.5. Math functions

marray tests will be added to every function from Table 175. Math functions which work on SYCL host and device.

2.3. 4.17.6. Integer functions

marray tests will be added to every function from Table 178. Integer functions which work on SYCL host and device, are available in the sycl namespace. Additionally some problems in existing tests should be fixed and new tests added:

  • Reference implementation of template <typename T, int N, typename R = typename std::make_unsigned_t<T>> sycl::vec<R, N> abs(sycl::vec<T, N> a) should return T, not R following Table 178 definition.

  • Reference implementation of template <typename T> T abs_diff(T a, T b) should return typename R = typename std::make_unsigned_t<T> instead following Table 178 definition. Implementation should be updated to avoid overflow and UB.

  • New function test for geninteger ctz(geninteger x) should be introduced along with reference implementation.

  • Reference implementations of mad24 and mul24 should be extended to cover full ranges for multipliers defined by specification.

2.4. 4.17.7. Common functions

marray tests will be added to every function from Table 179. Common functions which work on SYCL host and device, are available in the sycl namespace. Additionally some problems in existing tests should be fixed and new tests added:

  • Reference implementations of genfloat degrees(genfloat radians), genfloat radians(genfloat degrees), genfloat smoothstep(genfloat edge0, genfloat edge1, genfloat x), and genfloat sign(genfloat x) with sycl::half should be added.

  • Reference implementations of genfloat max(genfloat x, genfloat y) and genfloat min (genfloat x, genfloat y) should be fixed to provide correct undefined return value behavior (if any input is infinite or NaN).

  • Reference implementation of genfloat mix(genfloat x, genfloat y, genfloat a) with sycl::half should be added.

2.5. 4.17.8. Geometric functions

marray tests will be added to the following functions from Table 180. Geometric functions which work on SYCL host and device, are available in the sycl namespace:

Geometric Function Description
mfloat4 cross(mfloat4 p0, mfloat4 p1)
mfloat3 cross(mfloat3 p0, mfloat3 p1)
mdouble4 cross(mdouble4 p0, mdouble4 p1)
mdouble3 cross(mdouble3 p0, mdouble3 p1)

Returns the cross product of first 3 components of p0 and p1. The 4th component of result returned will be 0.0.

float dot(gengeofloat p0, gengeofloat p1)
double dot(gengeodouble p0, gengeodouble p1)

Compute dot product.

float distance(gengeofloat p0, gengeofloat p1)
double distance(gengeodouble p0, gengeodouble p1)

Returns the distance between p0 and p1. This is calculated as length(p0 - p1).

float length(gengeofloat p)
double length(gengeodouble p)

Return the length of vector p, i.e., \(\sqrt{ p.x^2 + p.y^2 + ...}\)

gengeofloat normalize(gengeofloat p)
gengeodouble normalize(gengeodouble p)

Returns a vector in the same direction as p but with a length of 1.

float fast_distance(gengeofloat p0, gengeofloat p1)

Returns fast_length(p0 - p1).

float fast_length(gengeofloat p)

Returns the length of vector p computed as: sqrt((half)(pow(p.x,2) + pow(p.y,2) + ...))

gengeofloat fast_normalize(gengeofloat p)

Returns a vector in the same direction as p but with a length of 1. fast_normalize is computed as:

p*rsqrt((half)(pow(p.x,2) + pow(p.y,2) + ... ))

The result shall be within 8192 ulps error from the infinitely precise result of

if (all(p == 0.0f))
  result = p;
else
  result = p / sqrt(pow(p.x, 2) + pow(p.y, 2) + ...);

with the following exceptions:

  1. If the sum of squares is greater than FLT_MAX then the value of the floating-point values in the result vector are undefined.

  2. If the sum of squares is less than FLT_MIN then the implementation may return back p.

  3. If the device is in “denorms are flushed to zero” mode, individual operand elements with magnitude less than sqrt(FLT_MIN) may be flushed to zero before proceeding with the calculation.

Additionally some problems in existing tests should be fixed and new tests added:

  • Reference implementation of template <typename T> T normalize(T p) should be fixed for scalar types to provide -1 for negative values.

  • Reference implementations of float fast_distance(gengeofloat p0, gengeofloat p1), float fast_length(gengeofloat p), and gengeofloat fast_normalize(gengeofloat p) should be fixed to follow spec definitions.

2.6. 4.17.9. Relational functions

The following realtional functions for sycl::half inputs should be provided with tests like already present ones for float and double:

Relational Function Description
vec<int16_t, { n }> isequal(half { n } x, half { n } y)

Returns the component-wise compare of x == y.

vec<int16_t, { n }> isnotequal(half { n } x, half { n } y)

Returns the component-wise compare of x != y.

vec<int16_t, { n }> isgreater(half { n } x, half { n } y)

Returns the component-wise compare of x > y.

vec<int16_t, { n }> isgreaterequal(half { n } x, half { n } y)

Returns the component-wise compare of x >= y.

vec<int16_t, { n }> isless(half { n } x, half { n } y)

Returns the component-wise compare of x < y.

vec<int16_t, { n }> islessequal(half { n } x, half { n } y)

Returns the component-wise compare of x <= y.

vec<int16_t, { n }> islessgreater(half { n } x, half { n } y)

Returns the component-wise compare of (x < y) || (x > y).

vec<int16_t, { n }> isfinite(half { n } x)

Test for finite value.

vec<int16_t, { n }> isinf(half { n } x)

Test for infinity value (positive or negative) .

vec<int16_t, { n }> isnan(half { n } x)

Test for a NaN.

vec<int16_t, { n }> isnormal(half { n } x)

Test for a normal value.

vec<int16_t, { n }> isordered(half { n } x, half { n } y)

Test if arguments are ordered. isordered() takes arguments x and y, and returns the result isequal(x, x) && isequal(y, y).

vec<int16_t, { n }> isunordered(half { n } x, half { n } y)

Test if arguments are unordered. isunordered() takes arguments x and y, returning non-zero if x or y is NaN, and zero otherwise.

vec<int16_t, { n }> signbit(half { n } x)

Test for sign bit. Returns the following for each component in x: -1 (i.e all bits set) if the sign bit in the component value is set else returns 0.

Reference implementations for the scalar input functions should be fixed to provide correct bool output, and marray tests should be added for the marray input functions of the following table:

Relational Function Description
bool isequal(sgenfloat x, sgenfloat y)
marray<bool, { N }> isequal(mgenfloat x, mgenfloat y)

Returns the component-wise compare of x == y.

bool isnotequal(sgenfloat x, sgenfloat y)
marray<bool, { N }> isnotequal(mgenfloat x, mgenfloat y)

Returns the component-wise compare of x != y.

bool isgreater(sgenfloat x, sgenfloat y)
marray<bool, { N }> isgreater(mgenfloat x, mgenfloat y)

Returns the component-wise compare of x > y.

bool isgreaterequal(sgenfloat x, sgenfloat y)
marray<bool, { N }> isgreaterequal(mgenfloat x, mgenfloat y)

Returns the component-wise compare of x >= y.

bool isless(sgenfloat x, sgenfloat y)
marray<bool, { N }> isless(mgenfloat x, mgenfloat y)

Returns the component-wise compare of x < y.

bool islessequal(sgenfloat x, sgenfloat y)
marray<bool, { N }> islessequal(mgenfloat x, mgenfloat y)

Returns the component-wise compare of x <= y.

bool islessgreater(sgenfloat x, sgenfloat y)
marray<bool, { N }> islessgreater(mgenfloat x, mgenfloat y)

Returns the component-wise compare of (x < y) || (x > y).

bool isfinite(sgenfloat x)
marray<bool, { N }> isfinite(mgenfloat x)

Test for finite value.

bool isinf(sgenfloat x)
marray<bool, { N }> isinf(mgenfloat x)

Test for infinity value (positive or negative) .

bool isnan(sgenfloat x)
marray<bool, { N }> isnan(mgenfloat x)

Test for a NaN.

bool isnormal(sgenfloat x)
marray<bool, { N }> isnormal(mgenfloat x)

Test for a normal value.

bool isordered(sgenfloat x, sgenfloat y)
marray<bool, { N }> isordered(mgenfloat x, mgenfloat y)

Test if arguments are ordered. isordered() takes arguments x and y, and returns the result isequal(x, x) && isequal(y, y).

bool isunordered(sgenfloat x, sgenfloat y)
marray<bool, { N }> isunordered(mgenfloat x, mgenfloat y)

Test if arguments are unordered. isunordered() takes arguments x and y, returning true if x or y is NaN, and false otherwise.

bool signbit(sgenfloat x)
marray<bool, { N }> signbit(mgenfloat x)

Test for sign bit, returning true if the sign bit in x is set, and false otherwise.

bool any(sigeninteger x)
bool any(migeninteger x)

Returns true if the most significant bit in any component of x is set; otherwise returns false.

bool all(sigeninteger x)
bool all(migeninteger x)

Returns true if the most significant bit in all components of x is set; otherwise returns false.

mgentype bitselect(mgentype a, mgentype b, mgentype c)

Each bit of the result is the corresponding bit of a if the corresponding bit of c is 0. Otherwise it is the corresponding bit of b.

mgentype select(mgentype a, mgentype b, marray<bool, { N }> c)

Returns the component-wise result = c ? b : a.