More question about shared mem Some point is unclear in the document

Accelerated Computing CUDA CUDA Programming and Performance
1

I want to ask about some points are unclear in CUDA document

  • In the CUDA application
    // Shared mem size is determined by the host app at run time
    extern shared int s_int;

There 's some things that unclear to me :

  • I think when smth declared as extern it should be defined some where. So where it is
  • It say the mem size is determined by the host, so how it is determined
  • how can it guarantee that there 's no conflict between shared memory
    For example if i have two shared memory array , can i define
    extern shared int a_int;
    extern shared int b_int;
    OR should i define the size of shared memory for each array
  • How can set the alignment for the shared memory
  • There a CUT_BANK_CHECKER(sdata, index) macro, can any one explain how to use this
    , just because i can see it a lot in the sample code
2

When you call you kernel:

MyKernel<<a,b,c>>();

c defines amount of shared memory allocated for your kernel (in bytes).

Next, when you declare variable:

extern __shared__ int s_int[];

s_int points to beginning of that allocated block. Any other declaration with extern specifier will also point to same location in shared memory.

BTW, this is RTFM question, it is described in Programming Manual. Have you read it before posting here?

Alignment can be enforced with specifier align(N).

3

So i guess there’s conflict if we declare 2 arrays like

_shared int a;

_shared int b;

that mean we can not have more than 1 shared mem array that is dynamically allocated is that right ?

Is that possible to define two function with the same name of shared mem with different types. Like this :

MyFunc1{

extern shared int sdata;

}

MyFunc1{

extern shared float sdata;

}

4

you can try it and you will know.
I think it is the same.
but It is very likely to make mistakes when using the shared data.

5

I believe a and b will point to the same memory. I have seen people successfully use code like this:

shared int a_int;
shared float a_float;

To create a “union” of sorts that allows them to treat the shared array as both types without explicit casts.

Be careful with this.

Mark