# accessing a multidimensional array in a kernel

**URL:** <https://forums.developer.nvidia.com/t/accessing-a-multidimensional-array-in-a-kernel/20717>\
**Category:** CUDA Programming and Performance\
**Created:** [January 26, 2011, 9:46pm UTC](https://forums.developer.nvidia.com/t/accessing-a-multidimensional-array-in-a-kernel/20717 "2011-01-26T21:46:42Z")\
**Posts on this page:** 13\
**Page:** 1

<div class="post-metadata">

**Author:** ![colede](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@colede](https://forums.developer.nvidia.com/u/colede)\
**Post date:** [January 26, 2011, 9:46pm UTC](https://forums.developer.nvidia.com/t/accessing-a-multidimensional-array-in-a-kernel/20717/1 "2011-01-26T21:46:42Z")

</div>

Hello,

If I use something like the following:

```
	cudaExtent extent = make_cudaExtent(W*sizeof(float), H, 9);

	cudaPitchedPtr covOut;
	cudaMalloc3D(&covOut, exent);//allocate the space on device

```

where W and H and 9 are the dimensions of my 3D array on the host,

can I access the variables in the 3D array on the device with the bracketed notation like

covOut[i][j][k]

Or do I have to use pointer arithmetic like

covOut[i\*W\*H+j\*W+k] to access it as if it were a single dimension array? If I have to access it like this, I thought I read somewhere that the CUDA arrays are stored in column major, so would that change the access to this:

covOut[j_W_H+i\*W+k]

Thanks!

---

<div class="post-metadata">

**Author:** ![avidday](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/avidday/32/59219_2.png) [@avidday](https://forums.developer.nvidia.com/u/avidday)\
**Post date:** [January 26, 2011, 10:06pm UTC](https://forums.developer.nvidia.com/t/accessing-a-multidimensional-array-in-a-kernel/20717/2 "2011-01-26T22:06:28Z")

</div>

Malloc2D and Malloc3D are only for allocating arrays for textures, you can’t use them to allocate general purpose global memory. Just use 3d indexing into linear memory using something like your second indexing scheme.

---

<div class="post-metadata">

**Author:** ![colede](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@colede](https://forums.developer.nvidia.com/u/colede)\
**Post date:** [January 26, 2011, 10:11pm UTC](https://forums.developer.nvidia.com/t/accessing-a-multidimensional-array-in-a-kernel/20717/3 "2011-01-26T22:11:22Z")

</div>

If this is the case, do I have to flatten my host 3D array into a 1D array before I can do a cudaMemCpy to the device 1D array, or is there a method or something for that?

---

<div class="post-metadata">

**Author:** ![avidday](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/avidday/32/59219_2.png) [@avidday](https://forums.developer.nvidia.com/u/avidday)\
**Post date:** [January 26, 2011, 10:21pm UTC](https://forums.developer.nvidia.com/t/accessing-a-multidimensional-array-in-a-kernel/20717/4 "2011-01-26T22:21:59Z")

</div>

Flattened arrays are considerably easier to work with, so flattening the array on the host is probably the best strategy.

---

<div class="post-metadata">

**Author:** ![colede](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@colede](https://forums.developer.nvidia.com/u/colede)\
**Post date:** [January 26, 2011, 10:26pm UTC](https://forums.developer.nvidia.com/t/accessing-a-multidimensional-array-in-a-kernel/20717/5 "2011-01-26T22:26:27Z")

</div>

> [@](#):
>
> Flattened arrays are considerably easier to work with, so flattening the array on the host is probably the best strategy.

I think I will go with this strategy for now.

I have one more question - the data that is being used in my kernel is actually using complex

is this going to be a problem within my cuda kernel?

---

<div class="post-metadata">

**Author:** ![colede](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@colede](https://forums.developer.nvidia.com/u/colede)\
**Post date:** [January 26, 2011, 10:36pm UTC](https://forums.developer.nvidia.com/t/accessing-a-multidimensional-array-in-a-kernel/20717/6 "2011-01-26T22:36:17Z")

</div>

> [@](#):
>
> I think I will go with this strategy for now.
> 
> I have one more question - the data that is being used in my kernel is actually using complex
> 
> is this going to be a problem within my cuda kernel?

Oh and one other thing - I guess this may be my C/C++ noobishness showing through -my code is currently C++ and i was just using new complex to allocate complex data. How would I do this in cudaMalloc?

---

<div class="post-metadata">

**Author:** ![Jimmy\_Pettersson](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Jimmy\_Pettersson](https://forums.developer.nvidia.com/u/Jimmy_Pettersson)\
**Post date:** [January 26, 2011, 10:45pm UTC](https://forums.developer.nvidia.com/t/accessing-a-multidimensional-array-in-a-kernel/20717/7 "2011-01-26T22:45:33Z")

</div>

You could use the float2 data type to represent your complex numbers.

So for example:

float2 complexVar;

complexVar.x = 1.0f;  
complexVar.y = 3.0f;

So complexVar can now be used to represent 1+3i. About cudaMalloc :

int size = N\*sizeof(float2); // number of complex elements times size of float2 (8 bytes)

cudaMalloc((void\*\*)&d\_complexArray, size);

Hope this is of some help!

---

<div class="post-metadata">

**Author:** ![colede](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@colede](https://forums.developer.nvidia.com/u/colede)\
**Post date:** [January 27, 2011, 12:12am UTC](https://forums.developer.nvidia.com/t/accessing-a-multidimensional-array-in-a-kernel/20717/8 "2011-01-27T00:12:38Z")

</div>

> [@](#):
>
> You could use the float2 data type to represent your complex numbers.
> 
> So for example:
> 
> float2 complexVar;
> 
> complexVar.x = 1.0f;
> 
> complexVar.y = 3.0f;
> 
> So complexVar can now be used to represent 1+3i. About cudaMalloc :
> 
> int size = N\*sizeof(float2); // number of complex elements times size of float2 (8 bytes)
> 
> cudaMalloc((void\*\*)&d\_complexArray, size);
> 
> Hope this is of some help!

What is float2? Have the operators been overridden to allow for things like complex addition and multiplication, absolute value, etc?

Also, what is cuComplex? I found some reference to it, and apparently I have this library in my lib directory. Can I use cuComplex for device and host code?

---

<div class="post-metadata">

**Author:** ![colede](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@colede](https://forums.developer.nvidia.com/u/colede)\
**Post date:** [January 27, 2011, 12:52am UTC](https://forums.developer.nvidia.com/t/accessing-a-multidimensional-array-in-a-kernel/20717/9 "2011-01-27T00:52:23Z")

</div>

I split this topic off to [The Official NVIDIA Forums | NVIDIA](http://forums.nvidia.com/index.php?showtopic=191730)

If a mod could move the relevant posts it’d be appreciated. Thanks

---

<div class="post-metadata">

**Author:** ![Jimmy\_Pettersson](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Jimmy\_Pettersson](https://forums.developer.nvidia.com/u/Jimmy_Pettersson)\
**Post date:** [January 27, 2011, 9:58am UTC](https://forums.developer.nvidia.com/t/accessing-a-multidimensional-array-in-a-kernel/20717/10 "2011-01-27T09:58:14Z")

</div>

> [@](#):
>
> What is float2? Have the operators been overridden to allow for things like complex addition and multiplication, absolute value, etc?
> 
> Also, what is cuComplex? I found some reference to it, and apparently I have this library in my lib directory. Can I use cuComplex for device and host code?

A float2 is 2 floats as the name indicates. Basically just:

struct FLOAT\_2 {

float x;

float y;

};

A cuComplex is the same thing as far as i know. float2 and cuComplex can be used in both host and device code.

You need to perform the complex multiplication, addition etc,. explicitly yourself, i.e.:

float2 x; x.x = 1.0f; x.y = 3.0f;

float2 y; y.x = 1.0f; y.y = 3.0f;

// x = x+y - complex addition

x.x = x.x + y.x;

x.y = x.y + y.y;

---

<div class="post-metadata">

**Author:** ![Dittoaway](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Dittoaway](https://forums.developer.nvidia.com/u/Dittoaway)\
**Post date:** [January 27, 2011, 3:06pm UTC](https://forums.developer.nvidia.com/t/accessing-a-multidimensional-array-in-a-kernel/20717/11 "2011-01-27T15:06:09Z")

</div>

.

---

<div class="post-metadata">

**Author:** ![avidday](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/avidday/32/59219_2.png) [@avidday](https://forums.developer.nvidia.com/u/avidday)\
**Post date:** [January 27, 2011, 3:08pm UTC](https://forums.developer.nvidia.com/t/accessing-a-multidimensional-array-in-a-kernel/20717/12 "2011-01-27T15:08:10Z")

</div>

Yeah, but CudaMath.h isn’t an offical part of CUDA, it is just a convenience file shipped in the SDK.

---

<div class="post-metadata">

**Author:** ![Dittoaway](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Dittoaway](https://forums.developer.nvidia.com/u/Dittoaway)\
**Post date:** [January 27, 2011, 3:53pm UTC](https://forums.developer.nvidia.com/t/accessing-a-multidimensional-array-in-a-kernel/20717/13 "2011-01-27T15:53:08Z")

</div>

Sorry, avidday, for leaving your post ‘hanging’.  
I realized exactly what you said and deleted my post … but too late ;}