Cuda 6.5 internal error while compiling

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,
while trying to compile a c++ template class using nvcc the following message is printed out:

thorsten@carme:/media/sda1/finroc$ nvcc -c -g -G -arch=sm_30 -dlink -std=c++11 -I sources/cpp/ -o build/linux_x86_64_debug/rrlib/cuda/math/cuVector.o sources/cpp/rrlib/cuda/math/cuVector.cu
sources/cpp/rrlib/cuda/math/cuVector.cuhpp(124): internal error: assertion failed at: “/dvs/p4/build/sw/rel/gpu_drv/r340/r340_00/drivers/compiler_CUREL/edg/EDG_4.8/src/scope_stk.c”, line 4704

1 catastrophic error detected in the compilation of “/tmp/tmpxft_00007c86_00000000-4_cuVector.cpp4.ii”.
Compilation aborted.

Does somebody know the problem or how to solve it?

Thank you,
Thorsten

2

Can you provide the cuVector.cu plus any necessary includes so that someone could duplicate the issue? Or a stripped-down version that still reproduces the issue?

Are you using the production version of CUDA 6.5 (i.e. not the RC version)?
Are you using the 340.29 driver that ships with the CUDA 6.5 production installer?
What is your OS?

3

I am using cuda-6-5 under ubuntu 14.04, 64 Bit, and installed it using the package from the website (production version).

Here the nvcc version:

nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2014 NVIDIA Corporation
Built on Thu_Jul_17_21:41:27_CDT_2014
Cuda compilation tools, release 6.5, V6.5.12

And here the glxinfo about the nvidia driver:

OpenGL vendor string: NVIDIA Corporation
OpenGL renderer string: GeForce GTX 750 Ti/PCIe/SSE2
OpenGL core profile version string: 4.3.0 NVIDIA 340.29

Here the cuVector.cu

#include "rrlib/cuda/math/cuVector.cuh"
const rrlib::cuda::math::cuVector<3, float> CVec3f;

cuVector.cuh:

#ifndef __rrlib__cuda__math__cuVector_h__
#define __rrlib__cuda__math__cuVector_h__

//----------------------------------------------------------------------
// External includes (system with <>, local with "")
//----------------------------------------------------------------------
#include <cuda_runtime.h>
#include <cuda.h>
#include <device_launch_parameters.h>
//----------------------------------------------------------------------
// Internal includes with ""
//----------------------------------------------------------------------

//----------------------------------------------------------------------
// Namespace declaration
//----------------------------------------------------------------------
namespace rrlib
{
namespace cuda
{
namespace math
{

//----------------------------------------------------------------------
// Forward declarations / typedefs / enums
//----------------------------------------------------------------------

//----------------------------------------------------------------------
// Class declaration
//----------------------------------------------------------------------
//! SHORT_DESCRIPTION
/*!
 * This class encapsulates a vector and provides some utility methods.
 */
template <unsigned int DIMENSION, typename tELEMENT>
class cuSpecialVectorMethods {
public:
	__device__ cuSpecialVectorMethods();
	__device__ ~cuSpecialVectorMethods();
};

template <unsigned int DIMENSION, typename tELEMENT>
class cuVector : public cuSpecialVectorMethods<DIMENSION, tELEMENT>
{

//----------------------------------------------------------------------
// Private fields and methods
//----------------------------------------------------------------------
private:
	tELEMENT data[DIMENSION];
	
	template<typename head, typename... tail>
	__device__ void InitHelper(const size_t index, head h, tail... t) {
		data[index] = h;
		InitHelper(index+1, t...);
	}
	
	__device__ void InitHelper(const size_t index) {}

//----------------------------------------------------------------------
// Public methods and typedefs
//----------------------------------------------------------------------
public:

//  __device__ cuVector();
  
  __device__ cuVector(const cuVector<DIMENSION, tELEMENT>& rhs);
  
  template<typename... Elements>
  __device__ cuVector(Elements... parameter);

  __device__ ~cuVector();
  
  template<typename T>
  __device__ tELEMENT& operator[](const T i);
  
    template<typename T>
  __device__ const tELEMENT& operator[](const T i) const;
  
  __device__ cuVector<DIMENSION, tELEMENT> DotProduct(const cuVector<DIMENSION, tELEMENT>& rhs) const;
  
  __device__ cuVector<DIMENSION, tELEMENT>& operator=(const cuVector<DIMENSION, tELEMENT>& rhs);
  
  __device__ cuVector<DIMENSION, tELEMENT> operator+(const cuVector<DIMENSION, tELEMENT>& rhs) const;
  
  __device__ cuVector<DIMENSION, tELEMENT>& operator+=(const cuVector<DIMENSION, tELEMENT>& rhs);
  
  __device__ cuVector<DIMENSION, tELEMENT> operator*(const cuVector<DIMENSION, tELEMENT>& rhs) const;
 
  __device__ cuVector<DIMENSION, tELEMENT>& operator*=(const cuVector<DIMENSION, tELEMENT>& rhs);
  
  __device__ unsigned int GetDimension() const;
  
  __device__ double GetLength() const;
};

template <typename tELEMENT>
class cuSpecialVectorMethods<3, tELEMENT> {
public:
	__device__ cuSpecialVectorMethods();
	__device__ ~cuSpecialVectorMethods();

	__device__ cuVector<3,tELEMENT> GetCrossProduct(const cuVector<3,tELEMENT>& rhs) const;
};


//----------------------------------------------------------------------
// End of namespace declaration
//----------------------------------------------------------------------
}
}
}

#include "rrlib/cuda/math/cuVector.cuhpp"

#endif

cuVector.cuhpp:

//----------------------------------------------------------------------
// External includes (system with <>, local with "")
//----------------------------------------------------------------------

//----------------------------------------------------------------------
// Internal includes with ""
//----------------------------------------------------------------------

//----------------------------------------------------------------------
// Debugging
//----------------------------------------------------------------------
#include <cassert>

//----------------------------------------------------------------------
// Namespace usage
//----------------------------------------------------------------------

//----------------------------------------------------------------------
// Namespace declaration
//----------------------------------------------------------------------
namespace rrlib
{
namespace cuda
{
namespace math
{

//----------------------------------------------------------------------
// Forward declarations / typedefs / enums
//----------------------------------------------------------------------

//----------------------------------------------------------------------
// Const values
//----------------------------------------------------------------------

//----------------------------------------------------------------------
// Implementation
//----------------------------------------------------------------------
template <unsigned int DIMENSION, typename tELEMENT>
__device__ cuSpecialVectorMethods<DIMENSION,tELEMENT>::cuSpecialVectorMethods() {}


template <unsigned int DIMENSION, typename tELEMENT>
__device__ cuSpecialVectorMethods<DIMENSION,tELEMENT>::~cuSpecialVectorMethods() {}


template <typename tELEMENT>
__device__ cuSpecialVectorMethods<3, tELEMENT>::cuSpecialVectorMethods() {}

template <typename tELEMENT>
__device__ cuSpecialVectorMethods<3, tELEMENT>::~cuSpecialVectorMethods() {}

template <typename tELEMENT>
__device__ cuVector<3,tELEMENT> cuSpecialVectorMethods<3, tELEMENT>::GetCrossProduct(const cuVector<3,tELEMENT>& rhs) const {
		cuVector<3,tELEMENT> result;
		const cuVector<3,tELEMENT>* current = static_cast<const cuVector<3,tELEMENT>*>(this);
		result[0] = current[1]*rhs[2]-current[2]*rhs[1];
		result[1] = current[2]*rhs[0]-current[0]*rhs[2];
		result[2] = current[0]*rhs[1]-current[1]*rhs[0];
		return result;
}

/*template <unsigned int DIMENSION, typename tELEMENT>
__device__ cuVector<DIMENSION,tELEMENT>::cuVector() {}*/
  
template <unsigned int DIMENSION, typename tELEMENT>
__device__ cuVector<DIMENSION,tELEMENT>::cuVector(const cuVector<DIMENSION, tELEMENT>& rhs) {
  	memcpy(data, rhs.data, DIMENSION*sizeof(tELEMENT));
}

template <unsigned int DIMENSION, typename tELEMENT>
template <typename... Elements>
__device__ cuVector<DIMENSION,tELEMENT>::cuVector(Elements... parameter){
	InitHelper(0, parameter...);
}

template <unsigned int DIMENSION, typename tELEMENT>
__device__ cuVector<DIMENSION,tELEMENT>::~cuVector() {}
  
template <unsigned int DIMENSION, typename tELEMENT>
template <typename T>
__device__ tELEMENT& cuVector<DIMENSION,tELEMENT>::operator[](const T i) {
  		return data[i];
 }
 
 template <unsigned int DIMENSION, typename tELEMENT>
template <typename T>
__device__ const tELEMENT& cuVector<DIMENSION,tELEMENT>::operator[](const T i) const {
  		return data[i];
 }
  
template <unsigned int DIMENSION, typename tELEMENT>
__device__ cuVector<DIMENSION, tELEMENT> cuVector<DIMENSION,tELEMENT>::DotProduct(const cuVector<DIMENSION, tELEMENT>& rhs) const {
  	cuVector<DIMENSION, tELEMENT> result;
  	for(usigned int i=0; i<DIMENSION; ++i) {
  		result[i] = data[i] * rhs[i];
  	}
  }
  
template <unsigned int DIMENSION, typename tELEMENT>
__device__ cuVector<DIMENSION, tELEMENT>& cuVector<DIMENSION,tELEMENT>::operator=(const cuVector<DIMENSION, tELEMENT>& rhs) {
  	memcpy(data, rhs.data, DIMENSION*sizeof(tELEMENT));
  	return *this;
  }
  
template <unsigned int DIMENSION, typename tELEMENT>
__device__ cuVector<DIMENSION, tELEMENT> cuVector<DIMENSION,tELEMENT>::operator+(const cuVector<DIMENSION, tELEMENT>& rhs) const {
  	cuVector<DIMENSION, tELEMENT> result;
  	for(unsigned int i=0; i<DIMENSION; ++i) {
  		result[i] = data[i] + rhs[i];
  	}
  	return result;	
  }
  
template <unsigned int DIMENSION, typename tELEMENT>
__device__ cuVector<DIMENSION, tELEMENT>& cuVector<DIMENSION,tELEMENT>::operator+=(const cuVector<DIMENSION, tELEMENT>& rhs) {
  	for(unsigned int i=0; i<DIMENSION; ++i) {
  		data[i] += rhs[i];
  	}
  	return *this;	
  }
  
template <unsigned int DIMENSION, typename tELEMENT>
__device__ cuVector<DIMENSION, tELEMENT> cuVector<DIMENSION,tELEMENT>::operator*(const cuVector<DIMENSION, tELEMENT>& rhs) const {
  	cuVector<DIMENSION, tELEMENT> result;
  	for(unsigned int i=0; i<DIMENSION; ++i) {
  		result[i] = data[i] * rhs[i];
  	}
  	return result;	
  }
  
template <unsigned int DIMENSION, typename tELEMENT>
__device__ cuVector<DIMENSION, tELEMENT>& cuVector<DIMENSION,tELEMENT>::operator*=(const cuVector<DIMENSION, tELEMENT>& rhs) {
  	tVector<DIMENSION, tELEMENT> result;
  	for(unsigned int i=0; i<DIMENSION; ++i) {
  		data[i] *= rhs[i];
  	}
  	return *this;	
  }
  
template <unsigned int DIMENSION, typename tELEMENT>
__device__ unsigned int cuVector<DIMENSION,tELEMENT>::GetDimension() const {
  	return DIMENSION;
  }
  
template <unsigned int DIMENSION, typename tELEMENT>
__device__ double cuVector<DIMENSION,tELEMENT>::GetLength() const {
  	double result = data[0];
  	for(unsigned int i=1; i<DIMENSION; ++i) {
  		result += data[i];
  	}
  	return sqrt(static_cast<double>(result));
  }
//----------------------------------------------------------------------
// End of namespace declaration
//----------------------------------------------------------------------
}
}
}
4

Hope this will help

5

Yes, thank you. I believe I am able to reproduce the error now. I don’t have any immediate comments, but if I learn or discover anything, I will report back.

6

Hi,

by now I have found the problem:
The compiler requires the constructor to be explicitly called, i.e. adding brackets:

cuVector<3,tELEMENT> result;

→ 

cuVector<3,tELEMENT> result();

Thanks for your fast response, help, and attention.

7

Sorry,
this should not have solved the problem.
By adding the brackets, the line of code showed above declares a function that returns a cuVector object instead of defining a variable of type cuVector.

8

Hi,

it was a simple error. For testing purposes I commented the standard constructor (cuVector.cuh:65) and forgot to uncomment it afterwards.
So it was my fault. Sorry.
But I wish the compiler had given a better error description.

Best Regards,
Thorsten