Template usage in CUDA code

Question

I am trying to understand how template is being used here:

template <uint32_t N_ATOMIC=32>
struct ParallelCounter {
public:
  uint32_t count[N_ATOMIC];

   // spread the counts across the counter
  __device__ __host__ void set(uint32_t x) {
    for(int i=0; i < N_ATOMIC; i++) count[i]=x/N_ATOMIC;

  } 

};



#ifndef SPREAD
#define SPREAD 32
#endif
__device__ ParallelCounter<SPREAD> myCounter;
__global__ void initCounter() {
  int tid = threadIdx.x + blockIdx.x * blockDim.x;

  if(tid == 0)
    myCounter.set(0);
}

All the examples on templates that I have seen so far have some variable in < > in first line above. But why we are having a constant uint32_t N_ATOMIC=32 here. if it is a constant, how would the template be useful? It is fixed for a particular type. Am I right? Thanks

The same as any defaults, regardless of cuda, templates, whatever — doctorlove, Sep 20 '13 at 14:19

score 1 · Answer 1 · edited May 23 '17 at 11:50

1

This question has nothing to do with CUDA.

If you define a template this way, it means that if you instantiate the template but don't provide that parameter specifically, it will default to the value indicated. You can still instantiate the template with other values, to get different behavior.

edited May 23 '17 at 11:50

Community

1
1

answered Sep 20 '13 at 14:18

Robert Crovella

143,785
11
213
257

kangshiyin · Answer 2 · 2013-09-20T14:35:45.540

It's not true that only typename can be used as a template parameter. A int const can also be a template parameter.

It's useful because this parameter can be determined in compile-time, rather than run-time. Earlier determination is helpful in generating more optimized run-time code.

In your particular case, the size of the struct ParallelCounter have to be determined during the compile-time. C++ won't accept a type whose size can not be determined until run-time.

If template is not used in your code, you may have to write a series of struct ParallelCounter and void initCounter() for different values of SPREAD

Template usage in CUDA code

2 Answers2