Some confusion on shared memory

Accelerated Computing CUDA CUDA Programming and Performance
1

I am reading through the PCR kernel code in cudpp and is confused by this section:

extern shared char shared;

T* a = (T*)shared;
T* b = (T*)&a[system_size];//
T* c = (T*)&b[system_size];
T* d = (T*)&c[system_size];
T* x = (T*)&d[system_size];

Why can’t we use

extern shared char shared;

T* a = (T*)shared;
T* b = (T*)shared;
T* c = (T*)shared;
T* d = (T*)shared;
T* x = (T*)shared;

And declared the shared memory size as the system_size in the host section?

2 
extern __shared__ char shared[];

	T* a = (T*)shared;

	T* b = (T*)shared;

	T* c = (T*)shared;

	T* d = (T*)shared;

	T* x = (T*)shared;

would alias a, b, c, d, x onto the same actual memory, while

extern __shared__ char shared[];

	T* a = (T*)shared;

	T* b = (T*)&a[system_size];//

	T* c = (T*)&b[system_size];

	T* d = (T*)&c[system_size];

	T* x = (T*)&d[system_size];

allocates different memory to each of them. The clever thing about it is that it will always allocate the right amount of memory for each of them (independently of their actual type), without explicitly calculating it (note how b is given in terms of a, c through b, …).

3

Forgive my ignorance, but could you spare a few words about the T* notation ? I am not
familiar with it. Is it something about C++ or PTX that I dont use ?

Thank you !

MW

4

It’s standard C. Just change the whitespace a bit to make it look more familiar:
[font=“Courier New”] T a = (T)shared;[/font]
This declares a pointer [font=“Courier New”]a[/font] that points to a variable of type [font=“Courier New”]T[/font], and initializes it to [font=“Courier New”]shared[/font], i.e. it points to the start of the shared memory block.

The spacing in the example is more C++ style. And more logical, because we initialize [font=“Courier New”]a[/font], not [font=“Courier New”]*a[/font].