Workarround to gemm interleaved with half type on Adaptivecpp

codeplaysoftware · Feb 27, 2024 · 8b4f189 · 8b4f189
1 parent f42fcd0
commit 8b4f189
Showing 1 changed file with 12 additions and 0 deletions.
diff --git a/src/operations/blas3/gemm_interleaved.hpp b/src/operations/blas3/gemm_interleaved.hpp
@@ -493,9 +493,21 @@ class Gemm<input_t, output_t, /* DoubleBuffer = */ false, /* NbcA = */ false,
       for (int j = 0; j < item_rows; ++j) {
 #pragma unroll
         for (int b = 0; b < item_batchs / VectorSize; ++b) {
+#ifdef __ADAPTIVECPP__
+          if constexpr (is_half<element_t>::value) {
+            *reg_res = reg_a[j * (item_batchs / VectorSize) + b] *
+                           reg_b[i * (item_batchs / VectorSize) + b] +
+                       *reg_res;
+          } else {
+            *reg_res = cl::sycl::mad(reg_a[j * (item_batchs / VectorSize) + b],
+                                     reg_b[i * (item_batchs / VectorSize) + b],
+                                     *reg_res);
+          }
+#else
           *reg_res = cl::sycl::mad(reg_a[j * (item_batchs / VectorSize) + b],
                                    reg_b[i * (item_batchs / VectorSize) + b],
                                    *reg_res);
+#endif
           ++reg_res;
         }
       }