How to declare a struct with a dynamic array inside it in device

Question

How to declare a struct in device that a member of it, is an array and then dynamically allocated memory for this. for example in below code, compiler said: error : calling a __host__ function("malloc") from a __global__ function("kernel_ScoreMatrix") is not allowed. is there another way for perform this action?

Type ofdev_size_idx_threads is int* and value of it, sent to kernel and used for allocate memory.

struct struct_matrix
{
    int *idx_threads_x;
    int *idx_threads_y;
    int thread_diag_length;
    int idx_length;
};

struct  struct_matrix matrix[BLOCK_SIZE_Y];

matrix->idx_threads_x= (int *) malloc ((*(dev_size_idx_threads) * sizeof(int) ));

score 4 · Accepted Answer · edited May 23 '17 at 11:50

From device code, dynamic memory allocations (malloc and new) are supported only with devices of cc2.0 and greater. If you have a cc2.0 device or greater, and you pass an appropriate flag to nvcc (such as -arch=sm_20) you should not see this error. Note that if you are passing multiple compilation targets (sm_10, sm_20, etc.), if even one of the targets does not meet the cc2.0+ requirement, you will see this error.

If you have a cc1.x device, you will need to perform these types of allocations from the host (e.g. using cudaMalloc) and pass appropriate pointers to your kernel.

If you choose that route (allocating from the host), you may also be interested in my answer to questions like this one.

EDIT: responding to questions below:

In visual studio (2008 express, should be similar for other versions), you can set the compilation target as follows: open project, select Project...Properties, select Configuration Properties...CUDA Runtime API...GPU Now, on the right hand pane, you will see entries like GPU Architecture (1) (and (2) etc.) These are drop-downs that you can click on and select the target(s) you want to compile for. If your GPU is sm_21 I would select that for (1) and leave the others blank, or select compatible versions like sm_20.
To see worked examples, please follow the link I gave above. A couple worked examples are linked from my answer here as well as a description of how it is done.

Thank you. my device CC is 2.1 , How to can set this parameter in Nsight visual studio? for allocating from the host, can you show a simple example? — mahdimb, May 29 '13 at 17:03
I found an answer for change `sm` parameter in [link](http://stackoverflow.com/questions/14411435/how-to-set-cuda-compiler-flags-in-visual-studio-2010) — mahdimb, May 29 '13 at 19:33

How to declare a struct with a dynamic array inside it in device

1 Answers1