Skip to content

Commit eb8fbf2

Browse files
author
Yihan Wang
authored
[SYCLomatic #1525] Add test for 14 asm instructions (#589)
Signed-off-by: Wang, Yihan <[email protected]>
1 parent 658adda commit eb8fbf2

File tree

4 files changed

+339
-11
lines changed

4 files changed

+339
-11
lines changed

‎features/feature_case/asm/asm_arith.cu

Lines changed: 146 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,7 @@
1616
#include <limits>
1717
#include <sstream>
1818
#include <string>
19+
#include <math.h>
1920

2021
#define CHECK(ID, S, CMP) \
2122
{ \
@@ -370,6 +371,141 @@ __device__ int shr() {
370371
return 0;
371372
}
372373

374+
template <typename T>
375+
__device__ T deg2rad(T val) {
376+
constexpr auto PI = 3.14159265358979323846f;
377+
return val * PI / 180.0f;
378+
}
379+
380+
#define FLOAT_CMP(X, Y) ((X - Y) < 1e-4)
381+
#define POWF2(X) (pow(2.0f, X))
382+
383+
__device__ int asm_copysign() {
384+
float f32 = 0.0f;
385+
double f64 = 0.0;
386+
CHECK(1, asm("copysign.f32 %0, %1, %2;" : "=f"(f32) : "f"(-10.0f), "f"(100.0f)), FLOAT_CMP(f32, -100.0f));
387+
CHECK(2, asm("copysign.f64 %0, %1, %2;" : "=d"(f64) : "d"(-10.0), "d"(100.0)), FLOAT_CMP(f64, -100.0));
388+
return 0;
389+
}
390+
391+
__device__ int asm_cos() {
392+
float f32 = 0.0f;
393+
CHECK(1, asm("cos.approx.f32 %0, %1;" : "=f"(f32) : "f"(deg2rad(90.0f))), FLOAT_CMP(f32, 0.0f));
394+
CHECK(2, asm("cos.approx.f32 %0, %1;" : "=f"(f32) : "f"(deg2rad(0.0f))), FLOAT_CMP(f32, 1.0f));
395+
CHECK(3, asm("cos.approx.f32 %0, %1;" : "=f"(f32) : "f"(deg2rad(60.0f))), FLOAT_CMP(f32, 0.5f));
396+
CHECK(4, asm("cos.approx.f32 %0, %1;" : "=f"(f32) : "f"(deg2rad(180.0f))), FLOAT_CMP(f32, -1.0f));
397+
CHECK(5, asm("cos.approx.f32.ftz %0, %1;" : "=f"(f32) : "f"(deg2rad(90.0f))), FLOAT_CMP(f32, 0.0f));
398+
CHECK(6, asm("cos.approx.f32.ftz %0, %1;" : "=f"(f32) : "f"(deg2rad(0.0f))), FLOAT_CMP(f32, 1.0f));
399+
CHECK(7, asm("cos.approx.f32.ftz %0, %1;" : "=f"(f32) : "f"(deg2rad(60.0f))), FLOAT_CMP(f32, 0.5f));
400+
CHECK(8, asm("cos.approx.f32.ftz %0, %1;" : "=f"(f32) : "f"(deg2rad(180.0f))), FLOAT_CMP(f32, -1.0f));
401+
return 0;
402+
}
403+
404+
__device__ int asm_sin() {
405+
float f32 = 0.0f;
406+
CHECK(1, asm("sin.approx.f32 %0, %1;" : "=f"(f32) : "f"(deg2rad(90.0f))), FLOAT_CMP(f32, 1.0f));
407+
CHECK(2, asm("sin.approx.f32 %0, %1;" : "=f"(f32) : "f"(deg2rad(0.0f))), FLOAT_CMP(f32, 0.0f));
408+
CHECK(3, asm("sin.approx.f32 %0, %1;" : "=f"(f32) : "f"(deg2rad(30.0f))), FLOAT_CMP(f32, 0.5f));
409+
CHECK(4, asm("sin.approx.f32 %0, %1;" : "=f"(f32) : "f"(deg2rad(180.0f))), FLOAT_CMP(f32, 0.0f));
410+
CHECK(5, asm("sin.approx.f32.ftz %0, %1;" : "=f"(f32) : "f"(deg2rad(90.0f))), FLOAT_CMP(f32, 1.0f));
411+
CHECK(6, asm("sin.approx.f32.ftz %0, %1;" : "=f"(f32) : "f"(deg2rad(0.0f))), FLOAT_CMP(f32, 0.0f));
412+
CHECK(7, asm("sin.approx.f32.ftz %0, %1;" : "=f"(f32) : "f"(deg2rad(30.0f))), FLOAT_CMP(f32, 0.5f));
413+
CHECK(8, asm("sin.approx.f32.ftz %0, %1;" : "=f"(f32) : "f"(deg2rad(180.0f))), FLOAT_CMP(f32, 0.0f));
414+
return 0;
415+
}
416+
417+
__device__ int asm_tanh() {
418+
float f32 = 0.0f;
419+
CHECK(1, asm("tanh.approx.f32 %0, %1;" : "=f"(f32) : "f"(deg2rad(45.0f))), FLOAT_CMP(f32, tanh(deg2rad(45.0f))));
420+
CHECK(2, asm("tanh.approx.f32 %0, %1;" : "=f"(f32) : "f"(deg2rad(0.0f))), FLOAT_CMP(f32, tanh(deg2rad(0.0f))));
421+
CHECK(3, asm("tanh.approx.f32 %0, %1;" : "=f"(f32) : "f"(deg2rad(180.0f))), FLOAT_CMP(f32, tanh(deg2rad(180.0f))));
422+
CHECK(4, asm("tanh.approx.f32 %0, %1;" : "=f"(f32) : "f"(deg2rad(90.0f))), FLOAT_CMP(f32, tanh(deg2rad(90.0f))));
423+
return 0;
424+
}
425+
426+
__device__ int asm_ex2() {
427+
float f32 = 0.0f;
428+
CHECK(1, asm("ex2.approx.f32 %0, %1;" : "=f"(f32) : "f"(2.1f)), FLOAT_CMP(f32, POWF2(2.1f)));
429+
CHECK(2, asm("ex2.approx.f32 %0, %1;" : "=f"(f32) : "f"(3.4f)), FLOAT_CMP(f32, POWF2(3.4f)));
430+
CHECK(3, asm("ex2.approx.f32 %0, %1;" : "=f"(f32) : "f"(9.7f)), FLOAT_CMP(f32, POWF2(9.7f)));
431+
CHECK(4, asm("ex2.approx.f32 %0, %1;" : "=f"(f32) : "f"(6.4f)), FLOAT_CMP(f32, POWF2(6.4f)));
432+
CHECK(5, asm("ex2.approx.f32.ftz %0, %1;" : "=f"(f32) : "f"(2.1f)), FLOAT_CMP(f32, POWF2(2.1f)));
433+
CHECK(6, asm("ex2.approx.f32.ftz %0, %1;" : "=f"(f32) : "f"(3.4f)), FLOAT_CMP(f32, POWF2(3.4f)));
434+
CHECK(7, asm("ex2.approx.f32.ftz %0, %1;" : "=f"(f32) : "f"(9.7f)), FLOAT_CMP(f32, POWF2(9.7f)));
435+
CHECK(8, asm("ex2.approx.f32.ftz %0, %1;" : "=f"(f32) : "f"(6.4f)), FLOAT_CMP(f32, POWF2(6.4f)));
436+
return 0;
437+
}
438+
439+
__device__ int asm_lg2() {
440+
float f32 = 0.0f;
441+
CHECK(1, asm("lg2.approx.f32 %0, %1;" : "=f"(f32) : "f"(2.1f)), FLOAT_CMP(f32, log2(2.1f)));
442+
CHECK(2, asm("lg2.approx.f32 %0, %1;" : "=f"(f32) : "f"(3.4f)), FLOAT_CMP(f32, log2(3.4f)));
443+
CHECK(3, asm("lg2.approx.f32 %0, %1;" : "=f"(f32) : "f"(9.7f)), FLOAT_CMP(f32, log2(9.7f)));
444+
CHECK(4, asm("lg2.approx.f32 %0, %1;" : "=f"(f32) : "f"(6.4f)), FLOAT_CMP(f32, log2(6.4f)));
445+
CHECK(5, asm("lg2.approx.f32.ftz %0, %1;" : "=f"(f32) : "f"(2.1f)), FLOAT_CMP(f32, log2(2.1f)));
446+
CHECK(6, asm("lg2.approx.f32.ftz %0, %1;" : "=f"(f32) : "f"(3.4f)), FLOAT_CMP(f32, log2(3.4f)));
447+
CHECK(7, asm("lg2.approx.f32.ftz %0, %1;" : "=f"(f32) : "f"(9.7f)), FLOAT_CMP(f32, log2(9.7f)));
448+
CHECK(8, asm("lg2.approx.f32.ftz %0, %1;" : "=f"(f32) : "f"(6.4f)), FLOAT_CMP(f32, log2(6.4f)));
449+
return 0;
450+
}
451+
452+
__device__ int sad() {
453+
int16_t s16;
454+
uint16_t u16;
455+
int32_t s32;
456+
uint32_t u32;
457+
int64_t s64;
458+
uint64_t u64;
459+
CHECK(1, asm("sad.s16 %0, %1, %2, %3;" : "=h"(s16) : "h"((int16_t)-1), "h"((int16_t)3), "h"((int16_t)5)), s16 == 9);
460+
CHECK(2, asm("sad.u16 %0, %1, %2, %3;" : "=h"(u16) : "h"((int16_t)1), "h"((int16_t)3), "h"((int16_t)5)), u16 == 7);
461+
CHECK(3, asm("sad.s32 %0, %1, %2, %3;" : "=r"(s32) : "r"(-1), "r"(3), "r"(5)), s32 == 9);
462+
CHECK(4, asm("sad.u32 %0, %1, %2, %3;" : "=r"(u32) : "r"(1), "r"(3), "r"(5)), u32 == 7);
463+
CHECK(5, asm("sad.s64 %0, %1, %2, %3;" : "=l"(s64) : "l"(-1ll), "l"(3ll), "l"(5ll)), s64 == 9);
464+
CHECK(6, asm("sad.u64 %0, %1, %2, %3;" : "=l"(u64) : "l"(1ll), "l"(3ll), "l"(5ll)), u64 == 7);
465+
return 0;
466+
}
467+
468+
__device__ int asm_rsqrt() {
469+
float f32;
470+
double f64;
471+
CHECK(1, asm("rsqrt.approx.f32 %0, %1;" : "=f"(f32) : "f"(2.1f)), FLOAT_CMP(f32, rsqrt(2.1f)));
472+
CHECK(2, asm("rsqrt.approx.f64 %0, %1;" : "=d"(f64) : "d"(2.1)), FLOAT_CMP(f64, rsqrt(2.1)));
473+
return 0;
474+
}
475+
476+
__device__ int asm_sqrt() {
477+
float f32;
478+
double f64;
479+
CHECK(1, asm("sqrt.approx.f32 %0, %1;" : "=f"(f32) : "f"(2.1f)), FLOAT_CMP(f32, sqrt(2.1f)));
480+
CHECK(2, asm("sqrt.approx.f32.ftz %0, %1;" : "=f"(f32) : "f"(2.1f)), FLOAT_CMP(f32, sqrt(2.1f)));
481+
CHECK(3, asm("sqrt.rn.f32 %0, %1;" : "=f"(f32) : "f"(2.1f)), FLOAT_CMP(f32, sqrt(2.1f)));
482+
CHECK(4, asm("sqrt.rz.f32 %0, %1;" : "=f"(f32) : "f"(2.1f)), FLOAT_CMP(f32, sqrt(2.1f)));
483+
CHECK(5, asm("sqrt.rm.f32 %0, %1;" : "=f"(f32) : "f"(2.1f)), FLOAT_CMP(f32, sqrt(2.1f)));
484+
CHECK(6, asm("sqrt.rp.f32 %0, %1;" : "=f"(f32) : "f"(2.1f)), FLOAT_CMP(f32, sqrt(2.1f)));
485+
CHECK(7, asm("sqrt.rn.f64 %0, %1;" : "=d"(f64) : "d"(2.1)), FLOAT_CMP(f64, sqrt(2.1)));
486+
CHECK(8, asm("sqrt.rz.f64 %0, %1;" : "=d"(f64) : "d"(2.1)), FLOAT_CMP(f64, sqrt(2.1)));
487+
CHECK(9, asm("sqrt.rm.f64 %0, %1;" : "=d"(f64) : "d"(2.1)), FLOAT_CMP(f64, sqrt(2.1)));
488+
CHECK(10, asm("sqrt.rp.f64 %0, %1;" : "=d"(f64) : "d"(2.1)), FLOAT_CMP(f64, sqrt(2.1)));
489+
return 0;
490+
}
491+
492+
__device__ int testp() {
493+
int pred = 0;
494+
{ asm(".reg .pred p1; testp.finite.f32 p1, %1; @p1 mov.s32 %0, 1;" : "=r"(pred) : "f"(0.1f)); if (!pred) { return 1; } };
495+
{ asm(".reg .pred p2; testp.infinite.f32 p2, %1; @p2 mov.s32 %0, 1;" : "=r"(pred) : "f"(std::numeric_limits<float>::infinity())); if (!pred) { return 2; } };
496+
{ asm(".reg .pred p3; testp.number.f32 p3, %1; @p3 mov.s32 %0, 1;" : "=r"(pred) : "f"(9.7f)); if (!pred) { return 3; } };
497+
{ asm(".reg .pred p4; testp.notanumber.f32 p4, %1; @p4 mov.s32 %0, 1;" : "=r"(pred) : "f"(NAN)); if (!pred) { return 4; } };
498+
{ asm(".reg .pred p5; testp.normal.f32 p5, %1; @p5 mov.s32 %0, 1;" : "=r"(pred) : "f"(9.5f)); if (!pred) { return 5; } };
499+
{ asm(".reg .pred p6; testp.subnormal.f32 p6, %1; @p6 mov.s32 %0, 1;" : "=r"(pred) : "f"(0.1e-300f)); if (!pred) { return 6; } };
500+
{ asm(".reg .pred p7; testp.finite.f64 p7, %1; @p7 mov.s32 %0, 1;" : "=r"(pred) : "d"(0.1)); if (!pred) { return 1; } };
501+
{ asm(".reg .pred p8; testp.infinite.f64 p8, %1; @p8 mov.s32 %0, 1;" : "=r"(pred) : "d"(std::numeric_limits<double>::infinity())); if (!pred) { return 2; } };
502+
{ asm(".reg .pred p9; testp.number.f64 p9, %1; @p9 mov.s32 %0, 1;" : "=r"(pred) : "d"(9.7)); if (!pred) { return 3; } };
503+
{ asm(".reg .pred p10; testp.notanumber.f64 p10, %1; @p10 mov.s32 %0, 1;" : "=r"(pred) : "d"(double(NAN))); if (!pred) { return 4; } };
504+
{ asm(".reg .pred p11; testp.normal.f64 p11, %1; @p11 mov.s32 %0, 1;" : "=r"(pred) : "d"(9.5)); if (!pred) { return 5; } };
505+
{ asm(".reg .pred p12; testp.subnormal.f64 p12, %1; @p12 mov.s32 %0, 1;" : "=r"(pred) : "d"(0.1e-400)); if (!pred) { return 6; } };
506+
return 0;
507+
}
508+
373509
__device__ int dp2a() {
374510
int32_t i32;
375511
uint32_t u32;
@@ -431,6 +567,16 @@ __global__ void test(int *ec) {
431567
TEST(cnot);
432568
TEST(shl);
433569
TEST(shr);
570+
TEST(asm_copysign);
571+
TEST(asm_cos);
572+
TEST(asm_sin);
573+
TEST(asm_tanh);
574+
TEST(asm_ex2);
575+
TEST(asm_lg2);
576+
TEST(sad);
577+
TEST(asm_rsqrt);
578+
TEST(asm_sqrt);
579+
TEST(testp);
434580
TEST(brev);
435581
TEST(dp2a);
436582
TEST(dp4a);

‎features/feature_case/asm/asm_v2inst.cu

Lines changed: 70 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -321,6 +321,61 @@ void testVmax2UUS(
321321
}
322322
}
323323

324+
__global__ void vavrg2UUS(int *Result, int a, int b, int c) {
325+
asm("vavrg2.u32.u32.s32 %0, %1, %2, %3;"
326+
: "=r"(*Result)
327+
: "r"(a), "r"(b), "r"(c));
328+
}
329+
330+
__global__ void vavrg2UUSSat(int *Result, int a, int b, int c) {
331+
asm("vavrg2.u32.u32.s32.sat %0, %1, %2, %3;"
332+
: "=r"(*Result)
333+
: "r"(a), "r"(b), "r"(c));
334+
}
335+
336+
__global__ void vavrg2UUSAdd(int *Result, int a, int b, int c) {
337+
asm("vavrg2.u32.u32.s32.add %0, %1, %2, %3;"
338+
: "=r"(*Result)
339+
: "r"(a), "r"(b), "r"(c));
340+
}
341+
342+
void testvavrg2UUS(
343+
const vector<pair<tuple<int, int, int>, tuple<int, int, int>>> &TestCases) {
344+
int *Result;
345+
cudaMallocManaged(&Result, sizeof(*Result));
346+
for (const auto &TestCase : TestCases) {
347+
string newCase = "{{" + to_string(get<0>(TestCase.first)) + ", " +
348+
to_string(get<1>(TestCase.first)) + ", " +
349+
to_string(get<2>(TestCase.first)) + "}, {";
350+
vavrg2UUS<<<1, 1>>>(Result, get<0>(TestCase.first), get<1>(TestCase.first),
351+
get<2>(TestCase.first));
352+
cudaDeviceSynchronize();
353+
newCase += to_string(*Result) + ", ";
354+
checkResult("vavrg2.u32.u32.s32",
355+
{get<0>(TestCase.first), get<1>(TestCase.first),
356+
get<2>(TestCase.first)},
357+
get<0>(TestCase.second), *Result);
358+
vavrg2UUSSat<<<1, 1>>>(Result, get<0>(TestCase.first),
359+
get<1>(TestCase.first), get<2>(TestCase.first));
360+
cudaDeviceSynchronize();
361+
newCase += to_string(*Result) + ", ";
362+
checkResult("vavrg2.u32.u32.s32.sat",
363+
{get<0>(TestCase.first), get<1>(TestCase.first),
364+
get<2>(TestCase.first)},
365+
get<1>(TestCase.second), *Result);
366+
vavrg2UUSAdd<<<1, 1>>>(Result, get<0>(TestCase.first),
367+
get<1>(TestCase.first), get<2>(TestCase.first));
368+
cudaDeviceSynchronize();
369+
newCase += to_string(*Result) + "}},";
370+
if (PRINT_CASE)
371+
cout << newCase << endl;
372+
checkResult("vavrg2.u32.u32.s32.add",
373+
{get<0>(TestCase.first), get<1>(TestCase.first),
374+
get<2>(TestCase.first)},
375+
get<2>(TestCase.second), *Result);
376+
}
377+
}
378+
324379
int main() {
325380
srand(unsigned(time(nullptr)));
326381
int a = rand();
@@ -401,6 +456,21 @@ int main() {
401456
{{2123447767, 63088206, 406272673}, {2123447767, 2123447767, 406320905}},
402457
{{1127203977, 209928516, 352777355}, {1127203977, 1127203977, 352844867}},
403458
});
459+
testvavrg2UUS({
460+
// {{a, b, c}, {0, 0, 0}},
461+
{{3, 4, 1}, {4, 4, 5}},
462+
{{30000000, 400000000, 10}, {214967232, 214967232, 12442}},
463+
{{INT16_MAX, 1, 100}, {16384, 16384, 16484}},
464+
{{UINT16_MAX, 1, 1000}, {32768, 32768, 33768}},
465+
{{UINT16_MAX, UINT16_MAX, 10000}, {32767, 32767, 42767}},
466+
{{INT16_MAX, UINT16_MAX, 100000}, {16383, 16383, 116383}},
467+
{{UINT32_MAX, UINT16_MAX, 1000000}, {-2147450881, -2147450881, 1065535}},
468+
{{INT16_MAX, UINT32_MAX, 10000000}, {-49153, 16383, 10016382}},
469+
{{INT16_MIN, INT32_MIN, 100000000}, {1073758208, 1073758208, 100032768}},
470+
{{454422046, 990649001, 541577428}, {722568292, 722568292, 541622345}},
471+
{{2123447767, 63088206, 406272673}, {1093333522, 1093271552, 406285789}},
472+
{{1127203977, 209928516, 352777355}, {668566247, 668566247, 352821067}},
473+
});
404474
cout << "passed " << passed << "/" << passed + failed << " cases!" << endl;
405475
if (failed) {
406476
cout << "failed!" << endl;

‎features/feature_case/asm/asm_v4inst.cu

Lines changed: 70 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -321,6 +321,61 @@ void testVmax4UUS(
321321
}
322322
}
323323

324+
__global__ void vavrg4UUS(int *Result, int a, int b, int c) {
325+
asm("vavrg4.u32.u32.s32 %0, %1, %2, %3;"
326+
: "=r"(*Result)
327+
: "r"(a), "r"(b), "r"(c));
328+
}
329+
330+
__global__ void vavrg4UUSSat(int *Result, int a, int b, int c) {
331+
asm("vavrg4.u32.u32.s32.sat %0, %1, %2, %3;"
332+
: "=r"(*Result)
333+
: "r"(a), "r"(b), "r"(c));
334+
}
335+
336+
__global__ void vavrg4UUSAdd(int *Result, int a, int b, int c) {
337+
asm("vavrg4.u32.u32.s32.add %0, %1, %2, %3;"
338+
: "=r"(*Result)
339+
: "r"(a), "r"(b), "r"(c));
340+
}
341+
342+
void testVavrg4UUS(
343+
const vector<pair<tuple<int, int, int>, tuple<int, int, int>>> &TestCases) {
344+
int *Result;
345+
cudaMallocManaged(&Result, sizeof(*Result));
346+
for (const auto &TestCase : TestCases) {
347+
string newCase = "{{" + to_string(get<0>(TestCase.first)) + ", " +
348+
to_string(get<1>(TestCase.first)) + ", " +
349+
to_string(get<2>(TestCase.first)) + "}, {";
350+
vavrg4UUS<<<1, 1>>>(Result, get<0>(TestCase.first), get<1>(TestCase.first),
351+
get<2>(TestCase.first));
352+
cudaDeviceSynchronize();
353+
newCase += to_string(*Result) + ", ";
354+
checkResult("vavrg4.u32.u32.s32",
355+
{get<0>(TestCase.first), get<1>(TestCase.first),
356+
get<2>(TestCase.first)},
357+
get<0>(TestCase.second), *Result);
358+
vavrg4UUSSat<<<1, 1>>>(Result, get<0>(TestCase.first),
359+
get<1>(TestCase.first), get<2>(TestCase.first));
360+
cudaDeviceSynchronize();
361+
newCase += to_string(*Result) + ", ";
362+
checkResult("vavrg4.u32.u32.s32.sat",
363+
{get<0>(TestCase.first), get<1>(TestCase.first),
364+
get<2>(TestCase.first)},
365+
get<1>(TestCase.second), *Result);
366+
vavrg4UUSAdd<<<1, 1>>>(Result, get<0>(TestCase.first),
367+
get<1>(TestCase.first), get<2>(TestCase.first));
368+
cudaDeviceSynchronize();
369+
newCase += to_string(*Result) + "}},";
370+
if (PRINT_CASE)
371+
cout << newCase << endl;
372+
checkResult("vavrg4.u32.u32.s32.add",
373+
{get<0>(TestCase.first), get<1>(TestCase.first),
374+
get<2>(TestCase.first)},
375+
get<2>(TestCase.second), *Result);
376+
}
377+
}
378+
324379
int main() {
325380
srand(unsigned(time(nullptr)));
326381
int a = rand();
@@ -401,6 +456,21 @@ int main() {
401456
{{2123447767, 63088206, 406272673}, {2123447767, 2123447767, 406273220}},
402457
{{1127203977, 209928516, 352777355}, {1127203977, 1127203977, 352777802}},
403458
});
459+
testVavrg4UUS({
460+
// {{a, b, c}, {0, 0, 0}},
461+
{{3, 4, 1}, {4, 4, 5}},
462+
{{30000000, 400000000, 10}, {206578752, 206578752, 202}},
463+
{{INT16_MAX, 1, 100}, {16512, 16512, 292}},
464+
{{UINT16_MAX, 1, 1000}, {32896, 32896, 1256}},
465+
{{UINT16_MAX, UINT16_MAX, 10000}, {32639, 32639, 10254}},
466+
{{INT16_MAX, UINT16_MAX, 100000}, {16255, 16255, 100190}},
467+
{{UINT32_MAX, UINT16_MAX, 1000000}, {-2139062401, -2139062401, 1000510}},
468+
{{INT16_MAX, UINT32_MAX, 10000000}, {-49281, 16255, 10000188}},
469+
{{INT16_MIN, INT32_MIN, 100000000}, {1082146816, 1082146816, 100000256}},
470+
{{454422046, 990649001, 541577428}, {722568419, 722568192, 541577591}},
471+
{{2123447767, 63088206, 406272673}, {1093333395, 1093271699, 406272912}},
472+
{{1127203977, 209928516, 352777355}, {685343591, 671122279, 352777590}},
473+
});
404474
cout << "passed " << passed << "/" << passed + failed << " cases!" << endl;
405475
if (failed) {
406476
cout << "failed!" << endl;

0 commit comments

Comments
 (0)