format

Signed-off-by: JackAKirk <[email protected]>

Author: JackAKirk

sycl/include/sycl/ext/oneapi/matrix/matrix-tensorcores.hpp

@@ -259,8 +259,7 @@ void load_multiplicand_cuda(
     sycl::ext::oneapi::experimental::matrix::joint_matrix<
         S, Use, NumRows, NumCols, Layout, sycl::sub_group> &res,
     multi_ptr<T, Space> src, size_t stride) {
-  if constexpr (std::is_same_v<
-                    S, sycl::ext::oneapi::experimental::bfloat16>) {
+  if constexpr (std::is_same_v<S, sycl::ext::oneapi::experimental::bfloat16>) {
     auto tileptr = reinterpret_cast<const int32_t *>(src.get());
     auto destptr = reinterpret_cast<int32_t *>(&res.wi_marray);
     if constexpr (NumRows == 16 && NumCols == 16) {
@@ -355,8 +354,8 @@ void load_multiplicand_cuda(
       __hmma_m32n8k16_ld_b(dstptr, tileptr, stride, get_layout_id<Layout>());
     }
 
-  } else if constexpr (std::is_same_v<S, sycl::ext::oneapi::experimental::matrix::
-                                           precision::tf32>) {
+  } else if constexpr (std::is_same_v<S, sycl::ext::oneapi::experimental::
+                                             matrix::precision::tf32>) {
     auto tileptr = reinterpret_cast<const int32_t *>(src.get());
     auto dstptr = reinterpret_cast<int32_t *>(&res.wi_marray);
     if constexpr (NumRows == 16 && NumCols == 8) {

sycl/test/check_device_code/matrix/matrix-nvptx-bfloat16-test.cpp

@@ -154,8 +154,8 @@ int main() {
           sycl::sub_group sg = item.get_sub_group();
 
           joint_matrix<float, use::accumulator, 8, 32> sub_c;
-joint_matrix<bfloat16, use::a, 8, 16, layout::col_major> sub_a;
-joint_matrix<bfloat16, use::b, 16, 32, layout::col_major> sub_b;
+          joint_matrix<bfloat16, use::a, 8, 16, layout::col_major> sub_a;
+          joint_matrix<bfloat16, use::b, 16, 32, layout::col_major> sub_b;
 
           // CHECK: tail call { float, float, float, float, float, float, float, float } @llvm.nvvm.wmma.m8n32k16.load.c.col.stride.f32.p1f32(float addrspace(1)* %_arg_accC, i32 16)
           joint_matrix_load(sg, sub_c, accC.get_pointer(), stride,

sycl/test/check_device_code/matrix/matrix-nvptx-uint8-test.cpp

@@ -42,15 +42,17 @@ int main() {
           joint_matrix<uint8_t, use::b, 16, 16, layout::row_major> sub_b;
 
           // CHECK: tail call { i32, i32, i32, i32, i32, i32, i32, i32 } @llvm.nvvm.wmma.m16n16k16.load.c.row.stride.s32.p1i32(i32 addrspace(1)* %_arg_accC, i32 16)
-          joint_matrix_load(sg, sub_c, accC.get_pointer(), stride, layout::row_major);
+          joint_matrix_load(sg, sub_c, accC.get_pointer(), stride,
+                            layout::row_major);
           // CHECK: tail call { i32, i32 } @llvm.nvvm.wmma.m16n16k16.load.a.row.stride.u8.p0i32(i32* %call.ascast.i.i{{.*}}.i, i32 16)
           joint_matrix_load(sg, sub_a, accA.get_pointer(), stride);
           // CHECK: tail call { i32, i32 } @llvm.nvvm.wmma.m16n16k16.load.b.row.stride.u8.p0i32(i32* %call.ascast.i.i{{.*}}.i, i32 16)
           joint_matrix_load(sg, sub_b, accB.get_pointer(), stride);
           // CHECK: tail call { i32, i32, i32, i32, i32, i32, i32, i32 } @llvm.nvvm.wmma.m16n16k16.mma.row.row.u8(i32 %11, i32 %12, i32 %15, i32 %16, i32 %1, i32 %2, i32 %3, i32 %4, i32 %5, i32 %6, i32 %7, i32 %8)
           sub_c = joint_matrix_mad(sg, sub_a, sub_b, sub_c);
           // CHECK: tail call void @llvm.nvvm.wmma.m16n16k16.store.d.row.stride.s32.p1i32(i32 addrspace(1)* %_arg_accD, i32 %18, i32 %19, i32 %20, i32 %21, i32 %22, i32 %23, i32 %24, i32 %25, i32 16)
-          joint_matrix_store(sg, sub_c, accD.get_pointer(), stride, layout::row_major);
+          joint_matrix_store(sg, sub_c, accD.get_pointer(), stride,
+                             layout::row_major);
         });
 
     cgh.parallel_for<class col_col_m16n16k16>(