Questions about usage of registers

[ArrogantGao]

Hello, I have some questions about usage of registers.
In GemmKernels.jl/kernel.jl file, function matmul_singlestage, I found

# (2) Load a compute_warp.M x compute_warp.N tile of C from shared memory into registers
warp_tile = @inbounds subdivide(block_tile.MN, Tile(conf.compute_warp).MN, warpId, conf.warps_per_block)

c_frags = LocalArray{Tuple{num_fragments_m, num_fragments_n}, Operator.fragtype_accum(conf.operator, conf.shared_c_layout)}(undef)

@loopinfo unroll for i = 1 : num_fragments_m
    @loopinfo unroll for j = 1 : num_fragments_n
        tile = translate_offset(warp_tile, (M = (i-1)*conf.compute_op_shape.M, N = (j-1)*conf.compute_op_shape.N))
        @inbounds @immutable c_frags[i, j] = transf_sh2rf_c(Operator.load_c(conf.operator, conf.shared_c_layout, shmem_c, tile), tile)
    end
end

which seems simply allocate some memory on the registers, but I found that the LocalArray is defined by

struct LocalArray{S <: Tuple, T, N, L} <: AbstractArray{T,N}
    data::NTuple{L,T}

    LocalArray{S,T,N,L}(::UndefInitializer) where {S,T,N,L} = new{S,T,N,L}()
    LocalArray{S,T,N,L}(data::NTuple{L,T}) where {S,T,N,L} = new{S,T,N,L}(data)
end

where the data are stored as NTuple.

So it that true that the NTuple type in gpu kernel are stored in registers? Or could you please tell me how to allocate memory on registers?

[maleadt]

So it that true that the NTuple type in gpu kernel are stored in registers? Or could you please tell me how to allocate memory on registers?

Generally, yes. In Julia, tuple-valued objects are emitted as LLVM IR that keeps the data stored in registers. It's possible that ptxas, the back-end compiler, will spill these to local memory (which is just device memory, i.e., slow), but that's rare.

[ArrogantGao]

Wow, that is really convenient. Thank you very much for your reply.

[maleadt]

Do note that there's other limits though, e.g. #99, where compiler heuristics prevent creation of large tuples (as they're generally pretty expensive, causing large IR, and for CPU execution are better moved to the heap).