why the --ptxas-options=-v doesn't show me local memory usage ?

Accelerated Computing CUDA CUDA Programming and Performance
1

I have a kernel which contains a statically allocated array of 50 bytes. I compiled it using nvcc.

nvcc -I../ -c -arch=sm_30 --ptxas-options=-v kernel.cu

It doesn’t show usage of local memory i.e. lmem

ptxas info    : 1024 bytes gmem
ptxas info    : Compiling entry function '_Z28function_S_' for 'sm_30'
ptxas info    : Function properties for _Z28function_S_
    56 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 22 registers, 368 bytes cmem[0], 4 bytes cmem[2]

I opened its ptx generated file and found the array declared in local memory

.local .align 1 .b8     __local_depot1[<u><b>50</b></u>];
        .reg .b64       %SP;
        .reg .b64       %SPL;
        .reg .pred      %p<14>;
        .reg .b16       %rs<25>;
        .reg .b32       %r<294>;

My question is; why the --ptxas-options=-v doesn’t show the lmem usage in its output.

$ nvcc --version

nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2015 NVIDIA Corporation
Built on Tue_Aug_11_14:27:32_CDT_2015
Cuda compilation tools, release 7.5, V7.5.17
2

ptxas is telling you the characteristics of SASS code, not PTX code.

Code analysis using PTX is rarely a good idea, IMO. It’s easy to reach the wrong conclusion.

Also, a compiler can easily choose to make size-known-at-compile-time local data stack-resident, instead of placed in the ordinary local logical space. Stack data is effectively similar to the logical local space in some respects.

It’s quite possible that the 50 byte local variable you see in PTX is being placed on the stack by ptxas, using 50 of the 56 bytes of the reported stack frame.

3

I have verified your statement and it seems correct.

There is one thing i want to know, When i set the array size to 20, the stack frame occupies 24 bytes. when i set it to 40, it occupies 40 bytes and for 50 bytes array, it occupies 56, Why the stack frame isn’t increasing linearly with the array size ?

4

Maybe there is an allocation granularity of 8 bytes.

5

Verified that. Thank you