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.
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