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?