How do I pass a multi-dimensional array by reference when it's a property of a struct?

Question

I have the following:

struct LR { double eps_dielect; 
        double sgm_conductivity; 
        double eno_ns_surfref;
        double frq_mhz; 
        double conf; 
        double rel;
        double erp;
        int radio_climate;  
        int pol;
        float antenna_pattern[361][1001];
          } LR;

I need to pass LR.antenna_pattern into a function, allocate the memory in a CUDA device then copy it. The float** type should represent LR.antenna_pattern[361][1001] just fine but I don't know how to instantiate the float** variable so that it is a pointer to LR.antenna_pattern

I try float** antennaPattern = (void**)&LR.antenna_pattern but it doesn't work. How do I create a pointer to LR.antenna_pattern?

Answer 1

One approach is to flatten your 2D array and handle it in a 1D fashion with pointer arithmetic to handle the row and column dimensions. First of all in your struct definition, replace the antenna_pattern element with:

struct LR { 
.
.
float *antenna_pattern;
} LR;

Then you will need to do a host-side malloc to allocate space:

#define COL 1001
#define ROW 361
#define DSIZE (ROW*COL)

LR.antenna_pattern = (float *)malloc(DSIZE*sizeof(float));

And a device side cuda malloc:

float *d_antenna_pattern;
cudaMalloc((void **) &d_antenna_pattern, DSIZE*sizeof(float));

The copy to the device looks like:

cudaMemcpy(d_antenna_pattern, LR.antenna_pattern, DSIZE*sizeof(float), cudaMemcpyHostToDevice);

When you want to reference into these arrays, you will have to do pointer arithmetic like:

float my_val_xy = ap[(x*COL)+y];  // to access element at [x][y] on the device

float my_val_xy = LR.antenna_pattern[(x*COL)+y]; // on the host

If you want to maintain the 2D array subscripts throughout, you can do this with an appropriate typedef. For an example, see the first code sample in my answer to this question. To diagram this out, you would need to start with a typedef:

#define COL 1001
#define ROW 361
#define DSIZE (ROW*COL)

typedef float aParray[COL];

and modify your structure definition:

struct LR { 
.
.
aParray *antenna_pattern;
} LR;

The host side malloc would look like:

LR.antenna_pattern = (aParray *)malloc(DSIZE*sizeof(float));

The device side cuda malloc would look like:

aParray *d_antenna_pattern;
cudaMalloc((void **) &d_antenna_pattern, DSIZE*sizeof(float));

The copy to the device looks like:

cudaMemcpy(d_antenna_pattern, LR.antenna_pattern, DSIZE*sizeof(float), cudaMemcpyHostToDevice);

The device kernel definition will need a function parameter like:

__global__ void myKernel(float ap[][COL]) {

Then inside the kernel you can access an element at x,y as:

float my_val_xy = ap[x][y];

Now in response to a follow-up question asking what to do if LR cannot be changed, here is a complete sample code which combines some of these ideas without modifying the LR structure:

#include<stdio.h>

// for cuda error checking
#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            return 1; \
        } \
    } while (0)



struct LR {
  int foo;
  float antenna_pattern[361][1001];
  } LR;

__global__ void mykernel(float ap[][1001]){

  int tid = threadIdx.x + (blockDim.x*blockIdx.x);
  float myval = 0.0;

  if (tid == 0){
    for (int i=0; i<361; i++)
      for (int j=0; j<1001; j++)
         ap[i][j] = myval++;
  }
}


int main(){

  typedef float aParray[1001];

  aParray *d_antenna_pattern;
  cudaMalloc((void **) &d_antenna_pattern, (361*1001)*sizeof(float));
  cudaCheckErrors("cudaMalloc fail");
  float *my_ap_ptr;
  my_ap_ptr = &(LR.antenna_pattern[0][0]);

  for (int i=0; i< 361; i++)
    for (int j=0; j<1001; j++)
      LR.antenna_pattern[i][j] = 0.0;
  cudaMemcpy(d_antenna_pattern, my_ap_ptr, (361*1001)*sizeof(float), cudaMemcpyHostToDevice);
  cudaCheckErrors("cudaMemcpy fail");
  mykernel<<<1,1>>>(d_antenna_pattern);
  cudaCheckErrors("Kernel fail");

  cudaMemcpy(my_ap_ptr, d_antenna_pattern, (361*1001)*sizeof(float), cudaMemcpyDeviceToHost);
  cudaCheckErrors("cudaMemcpy 2 fail");
  float myval = 0.0;
  for (int i=0; i<361; i++)
    for (int j=0; j<1001; j++)
      if (LR.antenna_pattern[i][j] != myval++) {printf("mismatch at offset x: %d y: %d actual: %f expected: %f\n", i, j, LR.antenna_pattern[i][j], --myval); return 1;}
  printf("Results match!\n");
  return 0;
}

If you prefer to use the flattened method, replace the d_antenna_pattern definition with:

float *d_antenna_pattern;

And change the kernel function parameter correspondingly to:

__global__ void mykernel(float *ap){

Then access using the pointer arithmetic method in the kernel:

ap[(i*1001)+j] = myval++;