Shouldn't this be 256?

I tried with 256 initially since CUDA used it. Reverted!

Did you check performance with both? If both of them give you correct results, select the one that gives you better performance.

Yes. A value of 64 instead of 256 so far gives the best performance on my A750.

Ideally, the best solution would required to test it on a B580 or on LunarLake iGPU, but I can't do it now. It's always possible to tweak it later. Thank you for testing and selecting the best.

I think it would be helpful if you are able to show us the results on a B580 or LL iGPU.

I tested on B580 and 64 gives the best result.

Thanks!

As this, in essence is an element-wise operation, we can have a 1 Dimensional kernel here. and the helper function I had mentioned will take care ascertaining the position along these dimensions.

Thanks. I'll try testing with a 1D kernel today.

Edit: It seems possible but do we really need a 1D kernel for scatter row shuffle/reorder kernel?

So the argument is that all (activations, conversions, etc etc...) our element-wise like kernels can be in condensed into just one 1 function for the most part as follows, as element-wise functions are just one-to-one mappings, be it anything.:

// kernel body
template<typename functor>
void elementwise(char* src,  char* dst, sycl::nd_item<1> it) {
    auto element_index = it.get_global_linear_id(); // id of the element in the flattened array.
    src_offset = calculate_offset(...) // using global element id, and dims and strides via that helper function
    auto val =reinterpret_cast<desired_type*>(src)[src_offset];
    auto output = functor(val);
    dst_offset = calculate_offset(...) // using global element id, and dims and strides, via that helper function
    // and similarly store in dst
}

The PR was merged before I attempt at this. I will open a new PR with 1D kerenl. (Almost same as cpy kernel performance)