StaticArrays memory operations with 1-byte alignment
- Dominant language
- Julia
- Stars
- 1.4k
- Forks
- 281
- Avg merge
- 1d 7h
- Merged PRs (30d)
- 30
Description
Julia performs most pointer operations with `align=1`, which performs badly on GPUs. We specialize `pointerref` and friends for CuDeviceArray, but that doesn't apply to StaticArrays (eg. `MArray`).
```
julia> using CUDAnative
using
julia> using StaticArrays
julia> kernel(a) = (@inbounds a[1] = zero(eltype(a)); nothing)
kernel (generic function with 1 method)
julia> CUDAnative.code_ptx(kernel, Tuple{CuDeviceArray{Float32,1,AS.Global}})
st.global.u32 [%rd3], %r1;
julia> CUDAnative.code_ptx(kernel, Tuple{MArray{Tuple{1}, Float32, 1, 1}})
st.u8 [%rd1+3], %r1;
st.u8 [%rd1+2], %r1;
st.u8 [%rd1+1], %r1;
st.u8 [%rd1], %r1;
```
Not sure how to best approach this. Keep `pointerref` and figure out alignment in an LLVM pass? Base should probably not emit memory operations with `align=1`.
Contributor guide
No contributing guide indexed for this repository
Assessment
This issue has not been assessed yet.