xe: concat: fix warnings and leading boundary shift

Author: syurkevi

src/gpu/intel/ocl/reusable_simple_concat.cl

@@ -474,8 +474,21 @@ internal_padding_block_concat2(__global DATA_T *dst,
         if (leading_boundary_shift) {
             int block_scaled_leading_shift
                     = leading_boundary_shift / blocks_per_simd1;
+            int rollover = leading_boundary_shift % blocks_per_simd1;
+
+            if (((get_local_id(0) / B0) + rollover) >= blocks_per_simd1)
+                block_scaled_leading_shift++;
+
             bVal = AS_VEC((as_ulong(bVal)
                     << block_scaled_leading_shift * DATA_TYPE_SIZE * 8));
+
+            // cannot directly shift values to corresponding location in sg since each
+            // workitem is responsible for multiple blocks, figure out source block
+            int src_bank = ((get_local_id(0) / B0)
+                                   + leading_boundary_shift % blocks_per_simd1)
+                    % blocks_per_simd1;
+            bVal = AS_VEC(intel_sub_group_shuffle(
+                    as_ulong(bVal), src_bank * B0 + (get_local_id(0) % B0)));
         }
 
         int sg_shuffle_dt = cutoff;

src/gpu/intel/ocl/reusable_simple_concat.cpp

@@ -236,7 +236,7 @@ static status_t attempt_normalize_ip_concat2(
         if (simd > total_elems) continue;
 
         const size_t min_bytes_per_workitem = 8;
-        const size_t elems_per_simd
+        const dim_t elems_per_simd
                 = simd * (min_bytes_per_workitem / data_type_size);
         const bool simd_even_block_multiple
                 = ((elems_per_simd) % conf.blocks[0]) == 0;
@@ -258,7 +258,7 @@ static status_t attempt_normalize_ip_concat2(
     conf.n = nonempty_inputs;
     conf.simd = max_simd;
     rt_conf.inner_axis = inner_offset;
-    conf.data_type_size = data_type_size;
+    conf.data_type_size = static_cast<int>(data_type_size);
 
     conf.use_large_index = (max_bytes > std::numeric_limits<int>::max());
 
@@ -296,7 +296,7 @@ static status_t attempt_normalize_ip_concat2(
 
     if (can_use_internal_padding_concat2) {
         rt_conf.inner_axis = concat2_inner_axis;
-        conf.data_type_size = concat2_dtsize;
+        conf.data_type_size = static_cast<int>(concat2_dtsize);
         conf.use_internal_padding_kernel = true;
 
         // TODO: compute::get_optimal_lws( // no emperical diff

src/gpu/intel/ocl/reusable_simple_concat.hpp

@@ -59,7 +59,7 @@ struct reusable_simple_concat_params_t
     int data_type_size;
     bool use_large_index = true;
     bool use_internal_padding_kernel = false;
-    uint8_t padding[3] = {0};
+    uint8_t padding[2] = {0};
 };
 
 struct reusable_simple_concat_runtime_params_t {