No __shared__ members of unions can anyone explain the rationale?

Accelerated Computing CUDA CUDA Programming and Performance
1

The shared qualifier is not permitted with members of unions and structures (CUDA programming guide section 4.2.2.4), and yet I have a need to create a union of shared arrays.

Ideally, I would like to be able to declare this in the .cu source file and successfully have two two-dimensional arrays with the same total number of elements but different dimensions:

#define BlockWidth …
#define BlockHeight …

union {

shared float As[BlockWidth][BlockHeight];
shared float Bs[BlockHeight][BlockWidth];

} Shared;

…

Shared.As[i][j] = …

Unfortunately, this is forbidden. I had hoped to use static multidimensional arrays rather than explicitly compute indices into a one dimensional array (i.e. As[i * BlockHeight + j] = …).

Can anyone explain:

  • why the compiler was written to prohibit shared members of unions in all cases
  • whether there is an alternative besides explicitly computing indices into a one-dimensional array

Thanks.

2

union/struct describe a type. shared is a modifier for a variable, so it doesn’t really make sense inside a type declaration.

Paulius

3

Agreed.

The solution is to make the instance of the union (or structure) shared.

shared union {

float As;

float Bs;

} Shared;

Shared.As = …

Shared.Bs = …

Thanks. Case closed.