From c8d4cf7bf9bbb8d4c59a39e9bd10f2b42ffb52a8 Mon Sep 17 00:00:00 2001 From: Vyacheslav N Klochkov Date: Thu, 25 Aug 2022 14:16:30 -0700 Subject: [PATCH 1/3] [ESIMD] Add a LIT test verifying DPAS with 2 tfloat32 arguments Signed-off-by: Vyacheslav N Klochkov --- SYCL/ESIMD/dpas/dpas_tf32.cpp | 87 +++++++++++++++++++++++++++++++++++ 1 file changed, 87 insertions(+) create mode 100644 SYCL/ESIMD/dpas/dpas_tf32.cpp diff --git a/SYCL/ESIMD/dpas/dpas_tf32.cpp b/SYCL/ESIMD/dpas/dpas_tf32.cpp new file mode 100644 index 0000000000..52c4d669b7 --- /dev/null +++ b/SYCL/ESIMD/dpas/dpas_tf32.cpp @@ -0,0 +1,87 @@ +//==---------------- dpas_tf32.cpp - DPC++ ESIMD on-device test ----------===// +// +// 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 +// +//===----------------------------------------------------------------------===// +// REQUIRES: gpu-intel-pvc || esimd_emulator +// UNSUPPORTED: cuda || hip +// RUN: %clangxx -fsycl -DESIMD_XE_HPC %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// XFAIL: esimd_emulator + +// The test verifies the low-level API for DPAS with 'tfloat32' types. +// It checks the versions of DPAS with and without the accumulator operand. + +#include +#include + +using namespace sycl; +using namespace sycl::ext::intel::esimd; +using namespace sycl::ext::intel::experimental::esimd; + +int main() { + queue Q; + + constexpr int REPEAT_COUNT = 8; + constexpr int SYSTOLIC_DEPTH = 8; + constexpr int EXECUTION_SIZE = 16; + + constexpr int M = REPEAT_COUNT; + constexpr int N = EXECUTION_SIZE; + constexpr int K = SYSTOLIC_DEPTH; // SYSTOLIC_DEPTH * OPS_PER_CHANNEL + float *A = malloc_shared(M * K, Q); + float *B = malloc_shared(K * N, Q); + float *C = malloc_shared(M * N, Q); + float *D = malloc_shared(M * N, Q); + for (int I = 0; I < M * K; ++I) + A[I] = I; + for (int I = 0; I < K * N; ++I) + B[I] = I; + + Q.single_task([=]() SYCL_ESIMD_KERNEL { + simd AVec(A); + simd BVec(B); + auto AView = AVec.template bit_cast_view(); + auto BView = BVec.template bit_cast_view(); + // C(MxN) = A(MxK) * B(KxN) + simd CVec = + dpas( + BView.read(), AView.read()); + CVec.copy_to(C); + + // D(MxN) = D(MxN) + A(MxK) * B(KxN); + simd DVec = 1.0; + DVec = dpas( + DVec, BView.read(), AView.read()); + DVec.copy_to(D); + }).wait(); + + unsigned ErrCnt = 0; + for (unsigned I = 0; I < M * N && ErrCnt < 10; ++I) { + int m = I / N; + int n = I % N; + float RefResC = 0.0f; + for (int k = 0; k < K; ++k) + RefResC += float((m * K + k) * (k * N + n)); + if (std::abs(RefResC - C[I]) > 0.001) { + std::cerr << "C[i] vs ref: " << C[I] << " : " << RefResC << std::endl; + ErrCnt++; + } + float RefResD = RefResC + 1.0; + if (std::abs(RefResD - D[I]) > 0.001) { + std::cerr << "D[i] vs ref: " << D[I] << " : " << RefResD << std::endl; + ErrCnt++; + } + } + free(A, Q); + free(B, Q); + free(C, Q); + free(D, Q); + + std::cout << (ErrCnt > 0 ? "FAILED\n" : "Passed\n"); + return ErrCnt > 0 ? 1 : 0; +} From 21db007d31884e169792e7c0d992e324202bb010 Mon Sep 17 00:00:00 2001 From: Vyacheslav N Klochkov Date: Fri, 26 Aug 2022 21:58:11 -0700 Subject: [PATCH 2/3] Added custom exception handler to the test Signed-off-by: Vyacheslav N Klochkov --- SYCL/ESIMD/dpas/dpas_tf32.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/SYCL/ESIMD/dpas/dpas_tf32.cpp b/SYCL/ESIMD/dpas/dpas_tf32.cpp index 52c4d669b7..db9c90de61 100644 --- a/SYCL/ESIMD/dpas/dpas_tf32.cpp +++ b/SYCL/ESIMD/dpas/dpas_tf32.cpp @@ -14,6 +14,8 @@ // The test verifies the low-level API for DPAS with 'tfloat32' types. // It checks the versions of DPAS with and without the accumulator operand. +#include "../esimd_test_utils.hpp" + #include #include @@ -22,7 +24,7 @@ using namespace sycl::ext::intel::esimd; using namespace sycl::ext::intel::experimental::esimd; int main() { - queue Q; + queue Q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler()); constexpr int REPEAT_COUNT = 8; constexpr int SYSTOLIC_DEPTH = 8; From 1667d589a1c3e03b280f28d9c6f789c208ad4913 Mon Sep 17 00:00:00 2001 From: Vyacheslav N Klochkov Date: Fri, 26 Aug 2022 22:00:59 -0700 Subject: [PATCH 3/3] Added braces to expression to guarantee no-warnings Signed-off-by: Vyacheslav N Klochkov --- SYCL/ESIMD/dpas/dpas_tf32.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/SYCL/ESIMD/dpas/dpas_tf32.cpp b/SYCL/ESIMD/dpas/dpas_tf32.cpp index db9c90de61..b9ee59a9d9 100644 --- a/SYCL/ESIMD/dpas/dpas_tf32.cpp +++ b/SYCL/ESIMD/dpas/dpas_tf32.cpp @@ -63,7 +63,7 @@ int main() { }).wait(); unsigned ErrCnt = 0; - for (unsigned I = 0; I < M * N && ErrCnt < 10; ++I) { + for (unsigned I = 0; (I < M * N) && (ErrCnt < 10); ++I) { int m = I / N; int n = I % N; float RefResC = 0.0f;