Namespace in CUDA

Hi everyone,

i’ve some trouble with CUDA, i’m trying to launch a __global __ function from a __device __ function called by an other __global __ function which is alo called by an other __global __ function and using namespace variables. (a lot…)

To explain, i defines some constant variable in a namespace instead of using #define.I get some trouble when the compilator said :
“error : expression must have a constant value”
I guess this error come from CUDA, 'cause my variable is const, and this code works properly in C++. If i change theses variable with random set, i get another error in the call of the kernel function :
“error : cannot pass an argument with a user-provided copy-constructor to a device-side kernel launch”

But there i don’t know where the error come from.

I need some help for those both problem.

Thank to all !

Want some code ?

namespace	global
{
	Global			global;
	const double	m_ITURadEarth = 6378145.0;
	const double	m_ITUGravCst = 3.986012E5;
	const double	m_ITUJ2 = 0.001082636;
	const double	m_J2000AngleDeg = 0; //-79.8058;
	const double	m_J2000AngleRad = 0; //TO_RAD(_J2000AngleDeg);
	const double	m_ITUAngleRateEarthRot = 4.1780745823E-3;
	const double	m_ITUAngleRateEarthRotRad = global.degToRad(m_ITUAngleRateEarthRot);
	const double	m_pi = 3.14159265358979323846;
	const double	m_2pi = 2 * m_pi;

	//////////////////////////////////////
	 Variables used by Kernel(s) function //
	//////////////////////////////////////
	// Main Kernel Function
	const double	m_simulationDuration = 615359.772; // time in second
	const double	m_step = 0.771;
	const double	m_nbItKernel = m_simulationDuration / m_step;
	const int		m_nbThreadsKernel = 1024;
	const int		m_nbBlocksKernel = (m_nbThreadsKernel - 1 + m_nbItKernel) / m_nbThreadsKernel;
	// Kernel for Satelite propagator
	const double	m_nbSat = 1;
	const int		m_nbThreadsSat = 128;
	const int		m_nbBlocksSat = (m_nbThreadsSat - 1 + m_nbSat) / m_nbThreadsSat;
	// Kernel for alpha angle (smallest angle possible on the Geostat orbit)
	const double	m_nbItAngle = 628319;
	const double	m_stepAngle = 1E-5;
	const int		m_nbThreadsAngle = 1024;
	const int		m_nbBlocksAngle = (m_nbThreadsAngle - 1 + m_nbItAngle) / m_nbThreadsAngle;
}


__device__ void				alphaAngle(Cartesian EarthStation, Cartesian Sat)
{
	double					CUDA_result[global::m_nbBlocksAngle], angle;
	int						CUDA_tid[m_nbBlocksAngle], index;
	Cartesian				GeoSat;

	/*GeoSat = findStartArcGeo(EarthStation);*/
	kernelAngle << < global::m_nbBlocksAngle, global::m_nbThreadsAngle >> > (EarthStation, Sat, GeoSat, CUDA_result, CUDA_tid);
	angle = getMinusAngle(CUDA_result);
	index = getIndexPosition(CUDA_result, CUDA_tid);
	 printf //
}

The error occurs on :

global::m_nbBlocksAngle

Oh and i wanna know something else. I’m coding on VisualStudio 2013 with CUDA7.5 and i wanna know if it’s possible for visual Studio to recognise “<<<” and _syncthreads(); as a valid syntax ?

I decided to pass all my variable from namespace to # define. This problem is solved. But i still have the other one.

When i try to launch a global function from a device function i get the error :

“error: cannot pass an argument with a user-provided copy-constructor to a device-side kernel launch”

Mmmmh… What does this mean ? i can’t launch my global function 'cause i create a Cartesian operator= ?

Same thing with the syntax : can “<<<” and _syncthreads be recognize as a valid syntax ?

I think it means exactly what it says. Since you’ve not provided a complete code, I won’t guess or speculate how it may apply to your code.

This is an intellisense issue, not recognizing CUDA constructs. It is widely discussed in this forum and numerous others, if you care to look. (for example, google “cuda red underline”) You can ignore it if you want, it should not prevent you from compiling or running proper code.

EDITED

No I don’t want to have to look through 600 lines of code if I don’t have to.

When asking for help on public forums, there are well-documented principles that suggest you should reduce your code down to something that is complete and minimal - contains something that someone else could compile, run and see the issue, but also only contains the minimum necessary content to reproduce the issue.

You already seem to have a grasp of the problem - it relates to an object you are passing as a kernel argument to a child kernel launch.

I think I could probably reproduce that error with 50 lines of code or less.

what is causing trouble is this :

class								Cartesian
{
public:
	double							m_X;
	double							m_Y;
	double							m_Z;
	Global							global;

private:
	double							m_m;

public:
	__host__ __device__				Cartesian(void);
	__host__ __device__				Cartesian(const Cartesian &c);
	__host__ __device__				Cartesian(double x, double y, double z);
	__host__ __device__	Cartesian&	operator=(const Cartesian& c);
	__host__ __device__ void		set(double x, double y, double z);
	__host__ __device__	double		magnitude(void);
	__host__ __device__ LatLonAlt	ToLLA(void);
};

__host__ __device__ Cartesian&	Cartesian::operator=(const Cartesian&c)
{
	m_X = c.m_X;
	m_Y = c.m_Y;
	m_Z = c.m_Z;
	m_m = c.m_m;
	global = c.global;
}

__host__ __device__ Cartesian		findStartArcGeo(LatLonAlt EarthStation)
{
	Global					global;
	double					deltaLon, LonRadMin, h;

	deltaLon = global.myAcos(cos(initeConstante(0)) / cos(EarthStation.m_lat));
	LonRadMin = global.ZeroTo2Pi(EarthStation.m_lon - deltaLon);
	h = ITU_RAD_EARTH + ITU_RAD_GSO;
	return (Cartesian(h * cos(LonRadMin), h * sin(LonRadMin), 0));
}

__host__ __device__ void	AlphaAngle(Cartesian EarthStation, Cartesian Sat)
{
	double			CUDA_result[BLOCKS_ANGLE], angle;
	int			CUDA_tid[BLOCKS_ANGLE], index;
	Cartesian		GeoSat = findStartArcGeo(EarthStation.ToLLA());

	kernelAngle <<< BLOCKS_ANGLE, THREADS_ANGLE >>> (EarthStation, Sat, GeoSat, CUDA_result, CUDA_tid);
	angle = getMinusAngle(CUDA_result);
	index = getIndexPosition(CUDA_result, CUDA_tid);
	// printf //
}

The line that show the error :
“error: cannot pass an argument with a user-provided copy-constructor to a device-side kernel launch”
is :

kernelAngle <<< BLOCKS_ANGLE, THREADS_ANGLE >>> (EarthStation, Sat, GeoSat, CUDA_result, CUDA_tid);

this error is repeated 3 time for the same line, guess is avout EarthStation, Sat and GeoSat

Function that i didn’t show arent special. Just return double, do some calculation…

I guess i know what is doing this (and maybe why) but i don’t know how to solve this

That’s a neat error, actually. Just reading it, I’m trying to imagine, if you have a user-defined copy constructor should you be able to pass that to a kernel? I guess it would all really depend on the implementation of the copy constructor.

I know that’s not super helpful but the error message is relatively straight-forward. Cannot pass an argument with a user-provided copy constructor to a device-side kernel launch. What a crazy language lol.

I understand the error-message. I get the fact that’s caus i made my own “operator=”. But if i don’t, i cannot do the stuff i want to do. And if i do, i cannot do the stuff i wan’t to do because i’ve got this error. So is there another way to do the “operator=” that allow me to launch a kernel function after an affectation ?

Or a kind of ‘trick’ that allow me to do it ?

Explaination :

without “operator=” i can’t do affectation 'cause i’ve got the error :

error : no operator “=” matches this operand

with the “operator=” i can’t launch my kernel function 'cause i’ve got the error :

error: cannot pass an argument with a user-provided copy-constructor to a device-side kernel launch

Si if i’ve understand well, my copy-constructor is wrong for kernel but good for my functions. So there is my constructor :

__host__ __device__ Cartesian&	Cartesian::operator=(const Cartesian&c)
{
	m_X = c.m_X;
	m_Y = c.m_Y;
	m_Z = c.m_Z;
	m_m = c.m_m;
	global = c.global;
}

Is there another way to do it that make the program work properly ?

Edit : Ok i put the operator= into comment and… The error is still here. I’m wondering if the probleme come from this…

If i declare a function in a device function must i allocate it with cudaMalloc ? Even if it’s just ‘for example’ an int ?

If the objects in question are in global memory, pass pointers to those objects to the device side kernel launch, rather than the objects thenmselves.

Yoo ! Thanks for that answer i was trying for almost an hour about how to make this thing work and i finally did. And i come over here to check and return what i get, and i find you’re answer x) So i finnally succeed. Just make all variable (class, struct) pointers and all will be ok ! Thanks to all