nvcc
Compiler error: __shared__ variables cannot have external linkage
extern __shared__ int col_ssd[1000]; //wrong
__shared__
array sizes in source code.
Has to be done at run time, eg via third argument on kernel launch
<<<
grid, block, bytes of shared memory per block>>>
.
Source code
extern __shared__ int col_ssd[];Notice empty
[]
.
extern __shared__ __align__(4) char shared_start[];For
sm_13
etc.,
use __align__(4)
to ensure that later on the
nvcc 6.0 compiler knows that shared_start
is aligned on a 4 byte boundary.
typedef unsigned int temptype_u4[4]; typedef char temptype_char5[5]; temptype_u4* ks = (temptype_u4*) &shared_start[shared_one]; temptype_u4* ls = (temptype_u4*) &ks[MAX_SEQUENCE_LENGTH]; temptype_char5* eligible_cs = (temptype_char5*) &ls[MAX_SEQUENCE_LENGTH];
shared_one
is a macro allowing space for
other variables in __shared__
memory.
In this example, since shared_start
elements are a byte wide,
shared_one
must be a multiple of four.
MAX_SEQUENCE_LENGTH
is a macro
(which need not be a multiple of four).
Example of using two dimensional array in shared memory:
eligible_cs[current_stage][no_of_eligible_cs++] = 1;