AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

EXSWHTEC-354 - Implement additional tests for integer intrinsics #430

Browse Source 
Change-Id: Icd155c2d2121bd9d25e7ac6ece9599553956457c


[ROCm/hip-tests commit: 147601aff2]
This commit is contained in:
Mirza Halilčević
2023-12-28 14:45:18 +01:00
committed by Rakesh Roy
parent 0f46c63d41
commit 4a1caf16ea
3 changed files with 484 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/hipTestMain/config/config_amd_linux
+5
View File
@@ -141,6 +141,11 @@
"Unit_Device___hbgtu2_Accuracy_Positive",
"Unit_Device___hbleu2_Accuracy_Positive",
"Unit_Device___hbltu2_Accuracy_Positive",
"=== Below 4 tests are disable due to defect EXSWHTEC-355 ===",
"Unit_Device___hadd_Sanity_Positive",
"Unit_Device___uhadd_Sanity_Positive",
"Unit_Device___rhadd_Sanity_Positive",
"Unit_Device___urhadd_Sanity_Positive",
"=== Patch which removes the typetraits implementation from std namespace in hiprtc is reverted ===",
"Unit_hiprtc_stdheaders",
"Unit_hipGraphAddMemcpyNode_Negative_Parameters",
projects/hip-tests/catch/hipTestMain/config/config_amd_windows
+5
View File
@@ -466,6 +466,11 @@
"Unit_Device___hbgtu2_Accuracy_Positive",
"Unit_Device___hbleu2_Accuracy_Positive",
"Unit_Device___hbltu2_Accuracy_Positive",
"=== Below 4 tests are disable due to defect EXSWHTEC-355 ===",
"Unit_Device___hadd_Sanity_Positive",
"Unit_Device___uhadd_Sanity_Positive",
"Unit_Device___rhadd_Sanity_Positive",
"Unit_Device___urhadd_Sanity_Positive",
#endif
"End of json"
]
projects/hip-tests/catch/unit/math/integer_intrinsics.cc
+474
View File
@@ -317,4 +317,478 @@ TEST_CASE("Unit_Device___umul24_Sanity_Positive") {
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(y.ptr()[0] == x1 * x2);
}
__global__ void __funnelshift_l_kernel(unsigned int* y, unsigned int lo, unsigned int hi,
unsigned int shift) {
y[0] = __funnelshift_l(lo, hi, shift);
}
/**
* Test Description
* ------------------------
* - Sanity test for `__funnelshift_l(lo,hi,shift)`.
*
* Test source
* ------------------------
* - unit/math/integer_intrinsics.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___funnelshift_l_Sanity_Positive") {
LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
const unsigned int lo = 0xAAAAAAAA, hi = 0xBBBBBBBB;
const unsigned long long hi_lo = (static_cast<unsigned long long>(hi) << 32) | lo;
for (unsigned int shift = 0; shift < 64; ++shift) {
__funnelshift_l_kernel<<<1, 1>>>(y.ptr(), lo, hi, shift);
HIP_CHECK(hipDeviceSynchronize());
INFO("shift: " << shift);
REQUIRE(y.ptr()[0] == static_cast<unsigned int>((hi_lo << (shift & 31)) >> 32));
}
}
__global__ void __funnelshift_lc_kernel(unsigned int* y, unsigned int lo, unsigned int hi,
unsigned int shift) {
y[0] = __funnelshift_lc(lo, hi, shift);
}
/**
* Test Description
* ------------------------
* - Sanity test for `__funnelshift_lc(lo,hi,shift)`.
*
* Test source
* ------------------------
* - unit/math/integer_intrinsics.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___funnelshift_lc_Sanity_Positive") {
LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
const unsigned int lo = 0xAAAAAAAA, hi = 0xBBBBBBBB;
const unsigned long long hi_lo = (static_cast<unsigned long long>(hi) << 32) | lo;
for (unsigned int shift = 0; shift < 64; ++shift) {
__funnelshift_lc_kernel<<<1, 1>>>(y.ptr(), lo, hi, shift);
HIP_CHECK(hipDeviceSynchronize());
INFO("shift: " << shift);
REQUIRE(y.ptr()[0] == static_cast<unsigned int>((hi_lo << std::min(shift, 32u)) >> 32));
}
}
__global__ void __funnelshift_r_kernel(unsigned int* y, unsigned int lo, unsigned int hi,
unsigned int shift) {
y[0] = __funnelshift_r(lo, hi, shift);
}
/**
* Test Description
* ------------------------
* - Sanity test for `__funnelshift_r(lo,hi,shift)`.
*
* Test source
* ------------------------
* - unit/math/integer_intrinsics.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___funnelshift_r_Sanity_Positive") {
LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
const unsigned int lo = 0xAAAAAAAA, hi = 0xBBBBBBBB;
const unsigned long long hi_lo = (static_cast<unsigned long long>(hi) << 32) | lo;
for (unsigned int shift = 0; shift < 64; ++shift) {
__funnelshift_r_kernel<<<1, 1>>>(y.ptr(), lo, hi, shift);
HIP_CHECK(hipDeviceSynchronize());
INFO("shift: " << shift);
REQUIRE(y.ptr()[0] == static_cast<unsigned int>(hi_lo >> (shift & 31)));
}
}
__global__ void __funnelshift_rc_kernel(unsigned int* y, unsigned int lo, unsigned int hi,
unsigned int shift) {
y[0] = __funnelshift_rc(lo, hi, shift);
}
/**
* Test Description
* ------------------------
* - Sanity test for `__funnelshift_rc(lo,hi,shift)`.
*
* Test source
* ------------------------
* - unit/math/integer_intrinsics.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___funnelshift_rc_Sanity_Positive") {
LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
const unsigned int lo = 0xAAAAAAAA, hi = 0xBBBBBBBB;
const unsigned long long hi_lo = (static_cast<unsigned long long>(hi) << 32) | lo;
for (unsigned int shift = 0; shift < 64; ++shift) {
__funnelshift_rc_kernel<<<1, 1>>>(y.ptr(), lo, hi, shift);
HIP_CHECK(hipDeviceSynchronize());
INFO("shift: " << shift);
REQUIRE(y.ptr()[0] == static_cast<unsigned int>(hi_lo >> std::min(shift, 32u)));
}
}
__global__ void __hadd_kernel(int* y, int x1, int x2) { y[0] = __hadd(x1, x2); }
/**
* Test Description
* ------------------------
* - Sanity test for `__hadd(x,y)`.
*
* Test source
* ------------------------
* - unit/math/integer_intrinsics.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___hadd_Sanity_Positive") {
LinearAllocGuard<int> y(LinearAllocs::hipMallocManaged, sizeof(int));
int x1 = GENERATE(0, -42, 42, 0xFFFFFFFF);
int x2 = GENERATE(0, -42, 42, 0xFFFFFFFF);
__hadd_kernel<<<1, 1>>>(y.ptr(), x1, x2);
HIP_CHECK(hipDeviceSynchronize());
INFO("x1: " << x1);
INFO("x2: " << x2);
REQUIRE(y.ptr()[0] == static_cast<int>((static_cast<long long>(x1) + x2) >> 1));
}
__global__ void __uhadd_kernel(unsigned int* y, unsigned int x1, unsigned int x2) {
y[0] = __uhadd(x1, x2);
}
/**
* Test Description
* ------------------------
* - Sanity test for `__uhadd(x,y)`.
*
* Test source
* ------------------------
* - unit/math/integer_intrinsics.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___uhadd_Sanity_Positive") {
LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
unsigned int x1 = GENERATE(0, 42, 0xFFFFFFFF);
unsigned int x2 = GENERATE(0, 42, 0xFFFFFFFF);
__uhadd_kernel<<<1, 1>>>(y.ptr(), x1, x2);
HIP_CHECK(hipDeviceSynchronize());
INFO("x1: " << x1);
INFO("x2: " << x2);
REQUIRE(y.ptr()[0] == static_cast<unsigned int>((static_cast<unsigned long long>(x1) + x2) >> 1));
}
__global__ void __rhadd_kernel(int* y, int x1, int x2) { y[0] = __rhadd(x1, x2); }
/**
* Test Description
* ------------------------
* - Sanity test for `__rhadd(x,y)`.
*
* Test source
* ------------------------
* - unit/math/integer_intrinsics.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___rhadd_Sanity_Positive") {
LinearAllocGuard<int> y(LinearAllocs::hipMallocManaged, sizeof(int));
int x1 = GENERATE(0, -42, 42, 0xFFFFFFFF);
int x2 = GENERATE(0, -42, 42, 0xFFFFFFFF);
__rhadd_kernel<<<1, 1>>>(y.ptr(), x1, x2);
HIP_CHECK(hipDeviceSynchronize());
INFO("x1: " << x1);
INFO("x2: " << x2);
REQUIRE(y.ptr()[0] == static_cast<int>((static_cast<long long>(x1) + x2 + 1) >> 1));
}
__global__ void __urhadd_kernel(unsigned int* y, unsigned int x1, unsigned int x2) {
y[0] = __urhadd(x1, x2);
}
/**
* Test Description
* ------------------------
* - Sanity test for `__urhadd(x,y)`.
*
* Test source
* ------------------------
* - unit/math/integer_intrinsics.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___urhadd_Sanity_Positive") {
LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
unsigned int x1 = GENERATE(0, 42, 0xFFFFFFFF);
unsigned int x2 = GENERATE(0, 42, 0xFFFFFFFF);
__urhadd_kernel<<<1, 1>>>(y.ptr(), x1, x2);
HIP_CHECK(hipDeviceSynchronize());
INFO("x1: " << x1);
INFO("x2: " << x2);
REQUIRE(y.ptr()[0] ==
static_cast<unsigned int>((static_cast<unsigned long long>(x1) + x2 + 1) >> 1));
}
__global__ void __mulhi_kernel(int* y, int x1, int x2) { y[0] = __mulhi(x1, x2); }
/**
* Test Description
* ------------------------
* - Sanity test for `__mulhi(x,y)`.
*
* Test source
* ------------------------
* - unit/math/integer_intrinsics.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___mulhi_Sanity_Positive") {
LinearAllocGuard<int> y(LinearAllocs::hipMallocManaged, sizeof(int));
int x1 = GENERATE(0, -42, 42, 0xFFFFFFFF);
int x2 = GENERATE(0, -42, 42, 0xFFFFFFFF);
__mulhi_kernel<<<1, 1>>>(y.ptr(), x1, x2);
HIP_CHECK(hipDeviceSynchronize());
INFO("x1: " << x1);
INFO("x2: " << x2);
REQUIRE(y.ptr()[0] ==
static_cast<int>((static_cast<long long>(x1) * static_cast<long long>(x2)) >> 32));
}
__global__ void __umulhi_kernel(unsigned int* y, unsigned int x1, unsigned int x2) {
y[0] = __umulhi(x1, x2);
}
/**
* Test Description
* ------------------------
* - Sanity test for `__umulhi(x,y)`.
*
* Test source
* ------------------------
* - unit/math/integer_intrinsics.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___umulhi_Sanity_Positive") {
LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
unsigned int x1 = GENERATE(0, 42, 0xFFFFFFFF);
unsigned int x2 = GENERATE(0, 42, 0xFFFFFFFF);
__umulhi_kernel<<<1, 1>>>(y.ptr(), x1, x2);
HIP_CHECK(hipDeviceSynchronize());
INFO("x1: " << x1);
INFO("x2: " << x2);
REQUIRE(y.ptr()[0] ==
static_cast<unsigned int>((static_cast<unsigned long long>(x1) * x2) >> 32));
}
__global__ void __mul64hi_kernel(long long* y, long long x1, long long x2) {
y[0] = __mul64hi(x1, x2);
}
/**
* Test Description
* ------------------------
* - Sanity test for `__mul64hi(x,y)`.
*
* Test source
* ------------------------
* - unit/math/integer_intrinsics.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___mul64hi_Sanity_Positive") {
LinearAllocGuard<long long> y(LinearAllocs::hipMallocManaged, sizeof(long long));
long long x1 = GENERATE(0, -42, 42, 0xFFFFFFFF);
long long x2 = GENERATE(0, -42, 42, 0xFFFFFFFF);
__mul64hi_kernel<<<1, 1>>>(y.ptr(), x1, x2);
HIP_CHECK(hipDeviceSynchronize());
INFO("x1: " << x1);
INFO("x2: " << x2);
REQUIRE(
y.ptr()[0] ==
static_cast<long long>((static_cast<__int128_t>(x1) * static_cast<__int128_t>(x2)) >> 64));
}
__global__ void __umul64hi_kernel(unsigned long long* y, unsigned long long x1,
unsigned long long x2) {
y[0] = __umul64hi(x1, x2);
}
/**
* Test Description
* ------------------------
* - Sanity test for `__umul64hi(x,y)`.
*
* Test source
* ------------------------
* - unit/math/integer_intrinsics.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___umul64hi_Sanity_Positive") {
LinearAllocGuard<unsigned long long> y(LinearAllocs::hipMallocManaged,
sizeof(unsigned long long));
unsigned long long x1 = GENERATE(0, 42, 0xFFFFFFFF);
unsigned long long x2 = GENERATE(0, 42, 0xFFFFFFFF);
__umul64hi_kernel<<<1, 1>>>(y.ptr(), x1, x2);
HIP_CHECK(hipDeviceSynchronize());
INFO("x1: " << x1);
INFO("x2: " << x2);
REQUIRE(y.ptr()[0] ==
static_cast<unsigned long long>(
(static_cast<__uint128_t>(x1) * static_cast<__uint128_t>(x2)) >> 64));
}
__global__ void __sad_kernel(unsigned int* y, int x1, int x2, unsigned int x3) {
y[0] = __sad(x1, x2, x3);
}
/**
* Test Description
* ------------------------
* - Sanity test for `__sad(x,y,z)`.
*
* Test source
* ------------------------
* - unit/math/integer_intrinsics.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___sad_Sanity_Positive") {
LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
int x1 = GENERATE(0, -42, 42, 0xFFFFFFFF);
int x2 = GENERATE(0, -42, 42, 0xFFFFFFFF);
unsigned int x3 = GENERATE(0, 42, 0xFFFFFFFF);
__sad_kernel<<<1, 1>>>(y.ptr(), x1, x2, x3);
HIP_CHECK(hipDeviceSynchronize());
INFO("x1: " << x1);
INFO("x2: " << x2);
REQUIRE(y.ptr()[0] == (static_cast<unsigned int>(std::abs(x1 - x2)) + x3));
}
__global__ void __usad_kernel(unsigned int* y, unsigned int x1, unsigned int x2, unsigned int x3) {
y[0] = __usad(x1, x2, x3);
}
/**
* Test Description
* ------------------------
* - Sanity test for `__usad(x,y,z)`.
*
* Test source
* ------------------------
* - unit/math/integer_intrinsics.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___usad_Sanity_Positive") {
LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
unsigned int x1 = GENERATE(0, 42, 0xFFFFFFFF);
unsigned int x2 = GENERATE(0, 42, 0xFFFFFFFF);
unsigned int x3 = GENERATE(0, 42, 0xFFFFFFFF);
__usad_kernel<<<1, 1>>>(y.ptr(), x1, x2, x3);
HIP_CHECK(hipDeviceSynchronize());
INFO("x1: " << x1);
INFO("x2: " << x2);
REQUIRE(y.ptr()[0] ==
(static_cast<unsigned int>(
std::abs(static_cast<long long>(x1) - static_cast<long long>(x2))) +
x3));
}
__global__ void __byte_perm(unsigned int* y, unsigned int x1, unsigned int x2, unsigned int s) {
y[0] = __byte_perm(x1, x2, s);
}
/**
* Test Description
* ------------------------
* - Sanity test for `__byte_perm(x,y,s)`.
*
* Test source
* ------------------------
* - unit/math/integer_intrinsics.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Device___byte_perm_Sanity_Positive") {
LinearAllocGuard<unsigned int> y(LinearAllocs::hipMallocManaged, sizeof(unsigned int));
unsigned int bytes[] = {0x88, 0x99, 0xAA, 0xBB, 0xCC, 0xDD, 0xEE, 0xFF};
unsigned int x1 = (bytes[3] << 24) | (bytes[2] << 16) | (bytes[1] << 8) | bytes[0];
unsigned int x2 = (bytes[7] << 24) | (bytes[6] << 16) | (bytes[5] << 8) | bytes[4];
unsigned int s0 = GENERATE(0, 1);
unsigned int s1 = GENERATE(2, 3);
unsigned int s2 = GENERATE(4, 5);
unsigned int s3 = GENERATE(6, 7);
unsigned int s = (s3 << 12) | (s2 << 8) | (s1 << 4) | s0;
__byte_perm<<<1, 1>>>(y.ptr(), x1, x2, s);
HIP_CHECK(hipDeviceSynchronize());
unsigned int expected = (bytes[s3] << 24) | (bytes[s2] << 16) | (bytes[s1] << 8) | bytes[s0];
REQUIRE(y.ptr()[0] == expected);
}