Namespaces and shared structs.

namespace.cu

#define KERNEL_QUALIFIERS                   __global__
#define VOLATILE                            volatile

#define DEVICE_INTRINSIC_QUALIFIERS         __device__ __forceinline__

#define DEVICE_STATIC_FUNCTION_QUALIFIERS   static DEVICE_FUNCTION_QUALIFIERS
#define DEVICE_STATIC_INTRINSIC_QUALIFIERS  static DEVICE_INTRINSIC_QUALIFIERS

//
//
//

DEVICE_STATIC_INTRINSIC_QUALIFIERS
unsigned int laneId()
{
  unsigned int id;

  asm("mov.u32 %0, %%laneid;" : "=r"(id));

  return id;
}

//
//
//

namespace foo
{
  VOLATILE __shared__ struct
  {
    float v[32];
  } shared;


  KERNEL_QUALIFIERS
  void
  baz(const float* const vin, float* const vout)
  {
    shared.v[laneId()] = vin[threadIdx.x];

    __syncthreads();

    vout[threadIdx.x] = shared.v[laneId()^31];
  }
}

//
//
//

namespace bar
{
  VOLATILE __shared__ struct
  {
    double v[32];
  } shared;

  KERNEL_QUALIFIERS
  void
  baz(const double* const vin, double* const vout)
  {
    shared.v[laneId()] = vin[threadIdx.x];

    __syncthreads();

    vout[threadIdx.x] = shared.v[laneId()^31];
  }
}

//
//
//

Compiled with: nvcc -m 32 -arch sm_35 -Xptxas=-v,-abi=no -cubin namespace.cu

ptxas : warning : 'option -abi=no' might get deprecated in future
ptxas : info : 0 bytes gmem
ptxas : info : Compiling entry function '_ZN3foo3bazEPKfPf' for 'sm_35'
ptxas : info : Used 4 registers, 128 bytes smem, 328 bytes cmem[0]
ptxas : info : Compiling entry function '_ZN3bar3bazEPKdPd' for 'sm_35'
ptxas : info : Used 6 registers, 256 bytes smem, 328 bytes cmem[0]