NVCC and C++ question about C++ usage in CUDA


i got a question about CUDAs relationship to C++. the programming manual at page 22 says: only the C subset of C++ is supported. what does that exactly mean? first it seemed perfectly clear, but the following example made me think

class vec3f



  vec3f(float x=0.0f, float y=0.0f, float z=0.0f) :x_(x), y_(y), z_(z) {}

  vec3f(const vec3f& v) : x_(v.x_), y_(v.y_), z_(v.z_) {}

  vec3f operator+(vec3f& src) 

  {return vec3f(x_ + src.x_, y_ + src.y_, z_ + src.z_);}

 __device__ float x() {return x_;}

  __device__ float y() {return y_;}

  __device__ float z() {return z_;}


  float x_;

  float y_;

  float z_;


__global__ void test_kernel (float4* pDst, float r, float g, float b)


  unsigned int x = threadIdx.x + blockIdx.x * blockDim.x;

  unsigned int y = threadIdx.y + blockIdx.y * blockDim.y;

  unsigned int width = blockDim.x * gridDim.x;

  vec3f v1(r,g,b);

  vec3f v2(v1+vec3f(-1.0f,0,0));

 pDst[x + width* y] = make_float4(v2.x(),v2.y(),v2.z(),0);


it’s not following the C++ spec (i have to use device), but operator overloading and constructors are clearly not C. the generated device code seems reasonable and executes fine.

so are these features still somewhat experimental and is there a certain reason for not using them, or are these things just poorly documented? (contrary to what the manual says basic inheritance seems to work, too…)


Certainly operator overloading, function overloading and templates work well. Mark Harris has a great example of using templates to optimize a kernel in his SC2007 talk:

thanks so far,

so what about classes and class members in CUDA ? i’ve tested simple derivation of classes, compiler produces nice code, too. virtual doesn’t seem to work.
so what is officially supported? i don’t want to base my work on things that are not officially ready for use…
which parts of the language C++ have already been or will be implemented in NVCC for use when writing device code (like operator overloading or classes in general)?