Reduce: vectorize contiguous by-block loads

[shreyas-omkar]

Use aligned 128-bit SIMD loads for stride-1 by-block reductions of 4- and
8-byte primitive elements. Each thread folds a Vec{4} or Vec{2} load into
its local accumulator; unsupported layouts continue through the scalar kernel.

The fast path requires a 16-byte-aligned buffer and row starts, a contiguous
reduced dimension, and an extent that is safe for the kernel's UInt32 index
arithmetic. A shared launcher applies the same dispatch to every by-block call
site.

The original benchmarks measured speedups of up to 44% on CUDA, 19% on AMDGPU,
37% on POCL, and 19% on oneAPI.

Validation:

• CPU suite: 34,692/34,692
• CUDA reduction suite with the sibling CUDA checkout: 18,131/18,131
• OpenCL/POCL reduction suite with the sibling OpenCL checkout: 18,131/18,131
• CUDA PTX contains one ld.global.v4 for Float32 and one ld.global.v2 for
  Float64

[shreyas-omkar]

@christiangnrd and @maleadt Please review it :)

[christiangnrd]

Can you add a few tests that are guaranteed to trigger this path?

[shreyas-omkar]

Yes I'll add them.

[christiangnrd]

@shreyas-omkar I added valid_backend_eltypes in #113 to make it easier to remember to filter out eltypes by backend

[shreyas-omkar]

@christiangnrd . Thank you for adding it. I have used it in the PR.