modifiable lvalue error typedef to define 2D array

Accelerated Computing CUDA CUDA Programming and Performance
1

Im a newbie to Cuda and I am trying to use typedef to define a 2D array to make it easier to interpret the data. I have written a basic cuda program to test this out. I want to pass a 2D array from matlab and then pass a copy of that 2D array back.

typedef double twoDarray[9][2];

global void copyArray(const double *scat,
double *out1)
{

twoDarray *scat2D;
scat2D = (twoDarray *)scat;

twoDarray *out;
out = (twoDarray *)out1;

for (int i = 0; i < 9; i++)
{
for (int j = 0; j < 2; j++)
{
out[i][j] = scat2D[i][j]; // error: expression must be a modifiable lvalue
}
}

}

Can someone please tell me what i’m doing wrong? Is it that i’m trying to access a const array. This doesn’t make sense to me.
And how are you suppose to add code properly?

2

I’m not so sure what you are trying to do here, but there is at least a problem in the line you pointed:
Your “twoDarray” is “similar” (ie naturally castable) to a double**, right? Then “out” which is a “twoDarray*” is similar to a double***.
So “out[i][j]” is similar to a double*, and something tells me that this is not what you want to get. But ultimately, I just don’t get what you want to do…

3

for this i’m just trying to write a function that will take an 2D array and pass a copy of that 2D array back. I am trying to figure out how to use the typedef properly to make it easier to access mutliple-D arrays in CUDA. It should be noted that i’m calling this function from Matlab using the feval.

Your “twoDarray” is “similar” (ie naturally castable) to a double**, right? This is correct.
I want to be able to cast both CUDA input arrays as 2D arrays so i can easily interpret the correct indecies.

in matlab, i’ll have an array defined, scats (9x2). Then call Cuda and have it give me a copy of this array, out (9x2).

4

I’m by no mean a matlab (never used it) but for what I heard and according to this matlab stores multi dimensional arrays contiguously in memory and in a column-major order. So doing this kind of copy can simply be done by passing a reference to the first element of the array, and by copying the whole multi dimensional array as if it was a 1D one of size the multiple of the individual dimensions.
For example, your (9,2) 2D array in internally stored as a 18 long 1D array, starting form element (1,1) (very much like in Fortran actually).
Now, I (along with wikipedia) might be wrong, but this is at least worth trying.

5

Hello!
I have a similar situation, but I pass arrays to GPU from Python, not Matlab.

The code below compiles and works. It is absolutely meaningless, but it shows the principle.
The float aa[140] array is treated as 3D array float matr[5][7][4].

global void init_pcur(Cmcmcfit *mc) {

#define L 5
#define M 7
#define N 4
//int L=5, M=7, N=4;

typedef float mat_t[M][N];
int i, j, k;

/* Make arrays aa[140] and matr[5][7][4] overlay */
float aa[140];
mat_t *matr = (mat_t *) aa;

for (i=0; i < N; ++i)
for (j=0; j < N; ++j)
for (k=0; k < M; ++k)
matr[i][j][k] = ijk;

. . . . . . . . . .

However, for my purposes this is useless: I need to overlay arrays with the dimensions dependent on blockDim.x, gridDim.x, threadIdx.x, and blockIdx.x. The nvcc compiler does not let it. The error message is
error: this operator is not allowed in an integral constant expression

If I try to use just variables int L=5, M=7, N=4, then the message is
error: expression must have a constant value

By the way, my gcc 4.4.3 allows expressions like typedef float mat_t[M][N], where N and N are variables. This is EXTREMELY useful in some cases. The code below compiles and works in Linux (you are encouraged to try):

#include <assert.h>
#include <stdio.h>
//#define N 9

void compare(int *array, int N, int M)
{
typedef int mat_t[N][M];
int i, j, k;
mat_t *matrix = (mat_t )array; / 3D array [any][N][M] /
for (i=0; i < N; ++i)
for (j=0; j < N; ++j)
for (k=0; k < M; ++k)
{
assert( &matrix[i][j][k] == &array[(iN+j)M+k] );
assert( matrix[i][j][k] == array[(iN+j)*M+k]);
printf(“matrix[%d][%d][%d]=%d\n”, i, j, k, matrix[i][j][k]);
}
}

int main(void)
{
int i, N=9, M=5;
int foo[NNM];
for (i=0; i < NNM; ++i) foo[i] = (7*i+11)*3+8;
compare(foo, N, M);
return 0;
}

…

Does anybody know if it is in principle possible in CUDF C ?
Does NVIDIA plan to upgrade nvcc to include variable bounds in typedef float mat_t[M][N]?