JuliaGPU / JuliaGPU/CUDA.jl

StaticArrays memory operations with 1-byte alignment

Open
#59 0 comments 0 reactions 0 assignees View on GitHub
cuda kernels performance
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.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.