Add missing offset to if conditions in the hierarchical and scoped kernel matrix assembly kernels.

Author: breyerml

include/plssvm/backends/SYCL/kernel/cg_explicit/hierarchical/kernel_matrix_assembly.hpp

@@ -83,7 +83,7 @@ class device_kernel_assembly {
         ::sycl::private_memory<std::array<std::array<real_type, static_cast<std::size_t>(INTERNAL_BLOCK_SIZE)>, static_cast<std::size_t>(INTERNAL_BLOCK_SIZE)>, 2> temp{ group };
 
         // only calculate the upper triangular matrix -> can't use get_local_id() since all work-items in a work-group must progress further
-        if (group[1] >= group[0]) {
+        if (group[1] + grid_y_offset_ >= group[0] + grid_x_offset_) {
             // initialize private temp matrix to zero
             group.parallel_for_work_item([&](::sycl::h_item<2> idx) {
                 for (unsigned internal_i = 0; internal_i < INTERNAL_BLOCK_SIZE; ++internal_i) {

include/plssvm/backends/SYCL/kernel/cg_explicit/scoped/kernel_matrix_assembly.hpp

@@ -86,7 +86,7 @@ class device_kernel_assembly {
                                    ::sycl::require_private_mem<std::array<std::array<real_type, static_cast<std::size_t>(INTERNAL_BLOCK_SIZE)>, static_cast<std::size_t>(INTERNAL_BLOCK_SIZE)>>(),  // temp
                                    [&](auto &data_i_cache, auto &data_j_cache, auto &temp) {
                                        // only calculate the upper triangular matrix -> can't use get_local_id() since all work-items in a work-group must progress further
-                                       if (group[1] >= group[0]) {
+                                       if (group[1] + grid_y_offset_ >= group[0] + grid_x_offset_) {
                                            // initialize private temp matrix to zero
                                            ::sycl::distribute_items_and_wait(group, [&](::sycl::s_item<2> idx) {
                                                for (unsigned internal_i = 0; internal_i < INTERNAL_BLOCK_SIZE; ++internal_i) {

include/plssvm/backends/SYCL/kernel/cg_implicit/hierarchical/kernel_matrix_assembly_blas.hpp

@@ -126,7 +126,7 @@ class device_kernel_assembly_symm {
         // implicit group barrier
 
         // only calculate the upper triangular matrix -> can't use get_local_id() since all work-items in a work-group must progress further
-        if (group[1] >= group[0]) {
+        if (group[1] + grid_y_offset_ >= group[0] + grid_x_offset_) {
             //*************************************************************************//
             //                   inplace kernel matrix construction                    //
             //*************************************************************************//

include/plssvm/backends/SYCL/kernel/cg_implicit/scoped/kernel_matrix_assembly_blas.hpp

@@ -126,7 +126,7 @@ class device_kernel_assembly_symm {
                                        });
 
                                        // only calculate the upper triangular matrix -> can't use get_local_id() since all work-items in a work-group must progress further
-                                       if (group[1] >= group[0]) {
+                                       if (group[1] + grid_y_offset_ >= group[0] + grid_x_offset_) {
                                            //*************************************************************************//
                                            //                   inplace kernel matrix construction                    //
                                            //*************************************************************************//