Passing Member Pointer as Parameter Gives Bad Values
In my current CUDA propject: Zs and Signals are members of Net, and I get good values passing Net.Zs and Net.Signals to printf but I get -1.#QNAN0 when I pass Zs and Signals to printf inside this kernel subroutine:

[code] typedef struct rohanNetwork
{// everything frequently acessed needs to be on 64-byte boundaries, 128 bytes if possible
cuDoubleComplex Wt[MAXWEIGHTS];
cuDoubleComplex Signals[MAXNEURONS];
cuDoubleComplex Deltas[MAXNEURONS];
cuDoubleComplex Zs[MAXNEURONS];
...
} rohanNetwork;


__device__ __align__(16) struct rohanNetwork devNet;


__device__ void subkEvalSampleBetaMT(rohanContext& Ses, long s, rohanNetwork& Net, int o, cuDoubleComplex * Signals, cuDoubleComplex * Zs, cuDoubleComplex * Wt, cuDoubleComplex * XInputs, cuDoubleComplex * YEval, double * dYEval, double * dSqrErr)
{
...
Signals[Net.iNeuronOfst[LAY]+kindex] = CxActivateUT( Zs[Net.iNeuronOfst[LAY]+kindex]);
if(s==0)printf("k>%d phi>%f+%f, Z> %f+%f\n", kindex,
Net.Signals[Net.iNeuronOfst[LAY]+kindex].x, Net.Signals[Net.iNeuronOfst[LAY]+kindex].y,
Net.Zs[Net.iNeuronOfst[LAY]+kindex].x, Net.Zs[Net.iNeuronOfst[LAY]+kindex].y );
if(s==0)printf("k>%d phi>%f+%f, Z> %f+%f\n", kindex,
Signals[Net.iNeuronOfst[LAY]+kindex].x, Signals[Net.iNeuronOfst[LAY]+kindex].y,
Zs[Net.iNeuronOfst[LAY]+kindex].x, Zs[Net.iNeuronOfst[LAY]+kindex].y );
...
}


subkEvalSampleBetaMT( devSes, s, devNet, true, Signals, Zs, devNet.Wt, devLearn.gpuXInputs, devLearn.gpuYEval, devLearn.gpudYEval, devLearn.gpudSqrErr);
[/code]


When this code is run, those last two printf statements produce something like this:
[code] k>0 phi>1.000000+0.000000, Z> 1.000000+0.000000
k>1 phi>0.996192+0.087184, Z> 0.064755+0.005667
k>0 phi>-1.#QNAN0+-1.#QNAN0, Z> -1.#QNAN0+-1.#QNAN0
k>1 phi>-1.#QNAN0+-1.#QNAN0, Z> -1.#QNAN0+-1.#QNAN0
[/code]

How can I access the correct values both as a function parameter's member pointer and as a member pointer passed as a parameter? Where is the flaw in my indirection?
In my current CUDA propject: Zs and Signals are members of Net, and I get good values passing Net.Zs and Net.Signals to printf but I get -1.#QNAN0 when I pass Zs and Signals to printf inside this kernel subroutine:



typedef struct rohanNetwork

{// everything frequently acessed needs to be on 64-byte boundaries, 128 bytes if possible

cuDoubleComplex Wt[MAXWEIGHTS];

cuDoubleComplex Signals[MAXNEURONS];

cuDoubleComplex Deltas[MAXNEURONS];

cuDoubleComplex Zs[MAXNEURONS];

...

} rohanNetwork;





__device__ __align__(16) struct rohanNetwork devNet;





__device__ void subkEvalSampleBetaMT(rohanContext& Ses, long s, rohanNetwork& Net, int o, cuDoubleComplex * Signals, cuDoubleComplex * Zs, cuDoubleComplex * Wt, cuDoubleComplex * XInputs, cuDoubleComplex * YEval, double * dYEval, double * dSqrErr)

{

...

Signals[Net.iNeuronOfst[LAY]+kindex] = CxActivateUT( Zs[Net.iNeuronOfst[LAY]+kindex]);

if(s==0)printf("k>%d phi>%f+%f, Z> %f+%f\n", kindex,

Net.Signals[Net.iNeuronOfst[LAY]+kindex].x, Net.Signals[Net.iNeuronOfst[LAY]+kindex].y,

Net.Zs[Net.iNeuronOfst[LAY]+kindex].x, Net.Zs[Net.iNeuronOfst[LAY]+kindex].y );

if(s==0)printf("k>%d phi>%f+%f, Z> %f+%f\n", kindex,

Signals[Net.iNeuronOfst[LAY]+kindex].x, Signals[Net.iNeuronOfst[LAY]+kindex].y,

Zs[Net.iNeuronOfst[LAY]+kindex].x, Zs[Net.iNeuronOfst[LAY]+kindex].y );

...

}





subkEvalSampleBetaMT( devSes, s, devNet, true, Signals, Zs, devNet.Wt, devLearn.gpuXInputs, devLearn.gpuYEval, devLearn.gpudYEval, devLearn.gpudSqrErr);






When this code is run, those last two printf statements produce something like this:

k>0 phi>1.000000+0.000000, Z> 1.000000+0.000000

k>1 phi>0.996192+0.087184, Z> 0.064755+0.005667

k>0 phi>-1.#QNAN0+-1.#QNAN0, Z> -1.#QNAN0+-1.#QNAN0

k>1 phi>-1.#QNAN0+-1.#QNAN0, Z> -1.#QNAN0+-1.#QNAN0




How can I access the correct values both as a function parameter's member pointer and as a member pointer passed as a parameter? Where is the flaw in my indirection?

Home: Asus mobo Pentium 4 w/ GTS550 Ti

Work: Asus mobo Core2Duo w/ GTS450 x2

Work: HP Pavillion Elite w/ Tesla C2075

64-bit Windows 7 Professional

CUDA 4.2, NSight 2.2.0.12132, VS 2008 SP1

#1
Posted 05/03/2012 07:42 AM   
Scroll To Top