[bernardogv]

The lane argument of core::arch::amdgpu::ds_permute and ds_bpermute is passed straight to llvm.amdgcn.ds_permute / llvm.amdgcn.ds_bpermute, which take a byte address like the ds_permute_b32 / ds_bpermute_b32 instructions do. To address lane n you pass n * 4. The docs currently say "lane", which reads like a lane index, and passing one silently gives wrong results.

This PR only changes the doc comments:

• Both functions now say that lane is a byte address (n * 4 for lane n).
• ds_bpermute's doc said it returns the value "given to ds_permute". It's the value given to ds_bpermute itself.

Evidence, from a gfx1151 (RDNA 3.5, wave32) test. One wave, lane i passes 100 + i and targets lane (i + 1) % 32:

call	lanes 0..5 get	correct?
ds_permute(t, ..)	[131, 106, 110, 114, 118, 122]	no, hardware writes to lane t / 4
ds_permute(t * 4, ..)	[131, 100, 101, 102, 103, 104]	yes
ds_bpermute(t, ..)	[100, 100, 100, 101, 101, 101]	no, reads lane t / 4
ds_bpermute(t * 4, ..)	[101, 102, 103, 104, 105, 106]	yes

HIP's own __shfl (hip/amd_detail/amd_warp_functions.h) does the same: __builtin_amdgcn_ds_bpermute(index << 2, var).

I first hit this with a Rust warp reduction (#![no_std], amdgcn-amd-amdhsa). Passing a bare lane to ds_bpermute gave max |err| 0.23 against the reference until I switched to lane * 4.

Tracking issue: rust-lang/rust#149988

[rustbot]

Thanks for the pull request, and welcome! The Rust Project has assigned @adamgemmell (or someone else) to review your changes, you should hear from them (or someone else) within the next two weeks.

Why was this reviewer chosen?

The reviewer was selected based on:

• Owners of files modified in this PR: @Amanieu, @adamgemmell, @davidtwco, @folkertdev
• @Amanieu, @adamgemmell, @davidtwco, @folkertdev expanded to Amanieu, adamgemmell, davidtwco, folkertdev
• Random selection from Amanieu, adamgemmell, davidtwco, folkertdev