sharedmem.h File Reference

Shared memory declaration struct for templatized types. More...

Classes
struct	SharedMemory< T >
	Wrapper class for templatized dynamic shared memory arrays. More...

---

Detailed Description

Shared memory declaration struct for templatized types.

sharedmem.h

Because dynamically sized shared memory arrays are declared "extern" in CUDA, we can't templatize their types directly. To get around this, we declare a simple wrapper struct that will declare the extern array with a different name depending on the type. This avoids linker errors about multiple definitions.

To use dynamically allocated shared memory in a templatized __global__ or __device__ function, just replace code like this:

  template<class T>
  __global__ void
  foo( T* d_out, T* d_in) 
  {
      // Shared mem size is determined by the host app at run time
      extern __shared__  T sdata[];
      ...
      doStuff(sdata);
      ...
  }
 

With this

  template<class T>
  __global__ void
  foo( T* d_out, T* d_in) 
  {
      // Shared mem size is determined by the host app at run time
      SharedMemory<T> smem;
      T* sdata = smem.getPointer();
      ...
      doStuff(sdata);
      ...
  }