Dynamic array inside a struct on constant memory?


I have been trying to adapt the SDK “particles” example for my own work. I wish for the user to be able to select a number of crowds to include in the simulation and to be able to assign individual sets of parameters to each crowd. The “particles” SDK example uses constant memory to store its SimParams object on the GPU:

[codebox]constant SimParams s_params;[/codebox]

A SimParams variable: m_params is created on the host and its variables set and copied to s_params on the gpu using the following:

[codebox]void setParameters(SimParams *hostParams)


// copy parameters to constant memory

cutilSafeCall( cudaMemcpyToSymbol(s_params, hostParams, sizeof(SimParams)) );


I have tried to add a dynamic array to the SimParams struct so I could assign the size when the number of crowds in the simulation is chosen. I set it up something like this:

[codebox]struct CrowdParams {

int size;

int dim;

float tag;

float visc;

float restDensity;

float charMass;

float aggression;

float fightRad;

float k;

float4 startPos;

float4 avgPos;


// simulation parameters

struct SimParams {

float3 gravity;

float globalDamping;

float charRadius;

uint3 gridSize;

uint numCells;

float3 worldOrigin;

float3 cellSize;

float spring;

float damping;

float shear;

float attraction;

float boundaryDamping;

CrowdParams *crowds;


Where the crowds array gets allocated on the host like so:

[codebox]m_params.crowds = (CrowdParams *)malloc(numCrowds * sizeof(CrowdParams));[/codebox]

I am able to printf all of the correct values from my m_params variable but when I copy it to the gpu and try to access information from s_params in a kernel the kernel fails to run properly:

[codebox]float radius = s_params.crowds[tag1].fightRad;[/codebox]

Access of variables that aren’t in the dynamic array seem to work okay.

I presume this is not working because I haven’t allocated memory for the crowds array on the gpu and that the copy is not taking into account my allocation from the host. Is this possible? or is there a better way?