From 229204313ac2a2a5a3ef542031632dd6b4401de6 Mon Sep 17 00:00:00 2001 From: "Tang, Jiajun" Date: Fri, 7 Jun 2024 16:30:45 +0800 Subject: [PATCH] [SYCLomatic #730] Add 6 math API tests. Signed-off-by: Tang, Jiajun jiajun.tang@intel.com --- features/feature_case/math/math-int.cu | 196 +++++++++++++++++++++++++ features/features.xml | 1 + features/test_feature.py | 2 +- 3 files changed, 198 insertions(+), 1 deletion(-) create mode 100644 features/feature_case/math/math-int.cu diff --git a/features/feature_case/math/math-int.cu b/features/feature_case/math/math-int.cu new file mode 100644 index 000000000..5c09e092b --- /dev/null +++ b/features/feature_case/math/math-int.cu @@ -0,0 +1,196 @@ +// ===--------------- math-int.cu ---------- *- CUDA -* -------------------===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +// +// ===---------------------------------------------------------------------===// + +#include +#include + +using namespace std; + +int passed = 0; +int failed = 0; + +void check(bool IsPassed) { + if (IsPassed) { + cout << " ---- passed" << endl; + passed++; + } else { + cout << " ---- failed" << endl; + failed++; + } +} + +template void printFunc(const string &FuncName, const T &Input) { + cout << FuncName << "(" << Input << ") "; +} + +template +void printFunc(const string &FuncName, const pair &Input) { + cout << FuncName << "(" << Input.first << ", " << Input.second << ")"; +} + +template void checkResult(const T &Expect, const T &Result) { + cout << " = " << Result << " (expect " << Expect << ")"; + check(Result == Expect); +} + +// Integer Mathematical Functions + +__global__ void _llmax(long long *const Result, long long Input1, + long long Input2) { + *Result = llmax(Input1, Input2); +} + +void testLlmaxCases( + const vector, long long>> &TestCases) { + long long *Result; + cudaMallocManaged(&Result, sizeof(*Result)); + for (const auto &TestCase : TestCases) { + _llmax<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second); + cudaDeviceSynchronize(); + printFunc("llmax", TestCase.first); + checkResult(TestCase.second, *Result); + } +} + +__global__ void _llmin(long long *const Result, long long Input1, + long long Input2) { + *Result = llmin(Input1, Input2); +} + +void testLlminCases( + const vector, long long>> &TestCases) { + long long *Result; + cudaMallocManaged(&Result, sizeof(*Result)); + for (const auto &TestCase : TestCases) { + _llmin<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second); + cudaDeviceSynchronize(); + printFunc("llmin", TestCase.first); + checkResult(TestCase.second, *Result); + } +} + +__global__ void _ullmax(unsigned long long *const Result, + unsigned long long Input1, unsigned long long Input2) { + *Result = ullmax(Input1, Input2); +} + +void testUllmaxCases( + const vector, + unsigned long long>> &TestCases) { + unsigned long long *Result; + cudaMallocManaged(&Result, sizeof(*Result)); + for (const auto &TestCase : TestCases) { + _ullmax<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second); + cudaDeviceSynchronize(); + printFunc("ullmax", TestCase.first); + checkResult(TestCase.second, *Result); + } +} + +__global__ void _ullmin(unsigned long long *const Result, + unsigned long long Input1, unsigned long long Input2) { + *Result = ullmin(Input1, Input2); +} + +void testUllminCases( + const vector, + unsigned long long>> &TestCases) { + unsigned long long *Result; + cudaMallocManaged(&Result, sizeof(*Result)); + for (const auto &TestCase : TestCases) { + _ullmin<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second); + cudaDeviceSynchronize(); + printFunc("ullmin", TestCase.first); + checkResult(TestCase.second, *Result); + } +} + +__global__ void _umax(unsigned *const Result, unsigned Input1, + unsigned Input2) { + *Result = umax(Input1, Input2); +} + +void testUmaxCases( + const vector, unsigned>> &TestCases) { + unsigned *Result; + cudaMallocManaged(&Result, sizeof(*Result)); + for (const auto &TestCase : TestCases) { + _umax<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second); + cudaDeviceSynchronize(); + printFunc("umax", TestCase.first); + checkResult(TestCase.second, *Result); + } +} + +__global__ void _umin(unsigned *const Result, unsigned Input1, + unsigned Input2) { + *Result = umin(Input1, Input2); +} + +void testUminCases( + const vector, unsigned>> &TestCases) { + unsigned *Result; + cudaMallocManaged(&Result, sizeof(*Result)); + for (const auto &TestCase : TestCases) { + _umin<<<1, 1>>>(Result, TestCase.first.first, TestCase.first.second); + cudaDeviceSynchronize(); + printFunc("umin", TestCase.first); + checkResult(TestCase.second, *Result); + } +} + +int main() { + testLlmaxCases({ + {{1, 2}, 2}, + {{-1, -2}, -1}, + {{1, -2}, 1}, + {{-1, 2}, 2}, + {{45212221678, 221332142421}, 221332142421}, + }); + testLlminCases({ + {{1, 2}, 1}, + {{-1, -2}, -2}, + {{1, -2}, -2}, + {{-1, 2}, -1}, + {{45212221678, 221332142421}, 45212221678}, + }); + testUllmaxCases({ + {{1, 2}, 2}, + {{18446744073709551615, 18446744073709551614}, 18446744073709551615}, + {{1, 18446744073709551614}, 18446744073709551614}, + {{18446744073709551615, 2}, 18446744073709551615}, + {{45212221678, 221332142421}, 221332142421}, + }); + testUllminCases({ + {{1, 2}, 1}, + {{18446744073709551615, 18446744073709551614}, 18446744073709551614}, + {{1, 18446744073709551614}, 1}, + {{18446744073709551615, 2}, 2}, + {{45212221678, 221332142421}, 45212221678}, + }); + testUmaxCases({ + {{1, 2}, 2}, + {{4294967295, 4294967294}, 4294967295}, + {{1, 4294967294}, 4294967294}, + {{4294967295, 2}, 4294967295}, + {{2262548718, 2288810325}, 2288810325}, + }); + testUminCases({ + {{1, 2}, 1}, + {{4294967295, 4294967294}, 4294967294}, + {{1, 4294967294}, 1}, + {{4294967295, 2}, 2}, + {{2262548718, 2288810325}, 2262548718}, + }); + cout << "passed " << passed << "/" << passed + failed << " cases!" << endl; + if (failed) { + cout << "failed!" << endl; + } + return failed; +} diff --git a/features/features.xml b/features/features.xml index 971570894..8436e9ce6 100644 --- a/features/features.xml +++ b/features/features.xml @@ -147,6 +147,7 @@ + diff --git a/features/test_feature.py b/features/test_feature.py index 2ef527bf3..e65c3f993 100644 --- a/features/test_feature.py +++ b/features/test_feature.py @@ -36,7 +36,7 @@ 'math-ext-bf16-conv', 'math-ext-double', 'math-ext-float', 'math-ext-half', 'math-ext-half-after11', 'math-ext-half-conv', 'math-ext-half2', 'math-ext-half2-after11', 'math-ext-simd', 'cudnn-activation', 'cudnn-fill', 'cudnn-lrn', 'cudnn-memory', 'cudnn-pooling', 'cudnn-reorder', 'cudnn-scale', 'cudnn-softmax', 'cudnn-sum', 'math-funnelshift', 'thrust-sort_by_key', 'thrust-find', 'thrust-inner_product', 'thrust-reduce_by_key', - 'math-bf16-conv', 'math-emu-bf16-conv-double', 'math-ext-bf16-conv-double', 'math-half-conv', + 'math-bf16-conv', 'math-emu-bf16-conv-double', 'math-ext-bf16-conv-double', 'math-half-conv', 'math-int', 'math-bfloat16', 'libcu_atomic', 'test_shared_memory', 'cudnn-reduction', 'cudnn-binary', 'cudnn-bnp1', 'cudnn-bnp2', 'cudnn-bnp3', 'cudnn-normp1', 'cudnn-normp2', 'cudnn-normp3', 'cudnn-convp1', 'cudnn-convp2', 'cudnn-convp3', 'cudnn-convp4', 'cudnn-convp5', 'cudnn-convp6', 'cudnn-convp7', 'cudnn_mutilple_files', "cusparse_1", "cusparse_2", "cusparse_3", "cusparse_4", "cusparse_5", "cusparse_6", "cusparse_7", "cusparse_8", "cusparse_9",