[SYCL][E2E][JM] Add GELU kernel to the functional suite (#18127)

Author: dkhaldi

sycl/test-e2e/Matrix/Inputs/common.hpp

@@ -261,3 +261,9 @@ template <typename T, layout Layout> constexpr int vnni_factor() {
   static_assert(sizeof(T) <= 4 && "Unsupported type in vnni_factor().");
   return 4 / sizeof(T);
 }
+
+inline float gelu(float val) {
+  return val *
+         (0.5f + 0.5f * sycl::tanh(val * (0.7978845608028654f +
+                                          0.035677408136300125f * val * val)));
+}

sycl/test-e2e/Matrix/Inputs/joint_matrix_bf16_fill_k_cache_impl.hpp

@@ -329,6 +329,9 @@ double joint_matmul(TOperand *A, TOperand *B, TResult *C, queue &q, int i
 #else  // MANUAL_UNROLL
             for (unsigned int n = 0; n < NCache1 / TN; n++) {
 #endif // MANUAL_UNROLL
+#ifdef GELU
+              joint_matrix_apply(sg, tC[m][n], [=](float &x) { x = gelu(x); });
+#endif // GELU
 #ifdef OOB
               ext::intel::experimental::matrix::joint_matrix_store_checked(
                   sg, tC[m][n], pC, colsB, layout::row_major, rowsA, colsB,
@@ -387,9 +390,13 @@ void test(size_t matrix_size_input) {
   matrix_rand<T>(matrix_size, matrix_size, A, T(1));
   matrix_rand<T>(matrix_size, matrix_size, B, T(1));
 
-  matrix_multiply_ref<T, T, TResult, 1>(A, B, refC, matrix_size, matrix_size,
-                                        matrix_size);
-
+  matrix_multiply_ref<T, T, TResult, 1>(
+      A, B, refC, matrix_size, matrix_size, matrix_size
+#ifdef GELU
+      ,
+      false, false, false, [](float &x) { x = gelu(x); }
+#endif // GELU
+  );
 #ifdef VNNI
   T *vnniB = malloc_shared<T>(matrix_size * matrix_size, q);
   matrix_vnni<T>(matrix_size, matrix_size, B, vnniB, vnniFactor);

sycl/test-e2e/Matrix/joint_matrix_gelu.cpp

@@ -0,0 +1,21 @@
+//==------------- joint_matrix_gelu.cpp  - DPC++ joint_matrix---------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+// UNSUPPORTED: target-nvidia, target-amd
+// UNSUPPORTED-INTENDED: aspect-ext_intel_matrix isn't currently supported for
+// other triples
+// XFAIL: run-mode
+// XFAIL-TRACKER: CMPLRLLVM-66371
+// REQUIRES: aspect-ext_intel_matrix
+
+// RUN: %{build} -o %t.out -DGELU %fp-model-precise
+// RUN: %{run} %t.out
+
+// -ffp-model=precise is added to not depend on compiler defaults.
+
+#include "common.hpp"
+#include "joint_matrix_bf16_fill_k_cache_impl.hpp"