# Pointer to CudaMem Using allocated Memory in different Kernels

**URL:** <https://forums.developer.nvidia.com/t/pointer-to-cudamem-using-allocated-memory-in-different-kernels/13354>\
**Category:** CUDA Programming and Performance\
**Created:** [November 20, 2009, 9:00am UTC](https://forums.developer.nvidia.com/t/pointer-to-cudamem-using-allocated-memory-in-different-kernels/13354 "2009-11-20T09:00:57Z")\
**Posts on this page:** 10\
**Page:** 1

<div class="post-metadata">

**Author:** ![CUDA\_VST\_RT](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@CUDA\_VST\_RT](https://forums.developer.nvidia.com/u/CUDA_VST_RT)\
**Post date:** [November 20, 2009, 9:00am UTC](https://forums.developer.nvidia.com/t/pointer-to-cudamem-using-allocated-memory-in-different-kernels/13354/1 "2009-11-20T09:00:57Z")

</div>

Hi,

iam using 2 kernel, which are synchronized, the first one allocate the momory on gpu:

extern “C” void X(…, float \*\*a\_d)  
{  
size\_t MEMsize = width_height_sizeof(float);  
unsigned int ArraySize = width \* height;

```
// allocate array on Device
(cudaMalloc((void **) a_d, MEMsize));

```

…  
}  
this will be called from c code:  
X(…, &a\_d);

doing so i got the adress of the memory in a\_d within the c-code. and use it when calling the second kernel

X2(…, a\_d);

within this second kernel i would like to acces the data in the GPU memory. its seems to work in the Emulation Mode, but changing to CUDA mode, it doesnt work.

so is there something to take care about?  
to avoid this situation i could copy the data just again within the second kernel, but that sounds inefficient.

any suggestions?  
thanks

---

<div class="post-metadata">

**Author:** ![sagrailo](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@sagrailo](https://forums.developer.nvidia.com/u/sagrailo)\
**Post date:** [November 20, 2009, 9:18am UTC](https://forums.developer.nvidia.com/t/pointer-to-cudamem-using-allocated-memory-in-different-kernels/13354/2 "2009-11-20T09:18:27Z")

</div>

It doesn’t seem, from your code snippets, that you are actually using two kernels - what you described is a host function to allocate memory, and then a kernel call. The sequence of operations look OK, but - are you checking cudaMalloc() result for errors (maybe the amount of memory you are requesting there is too big)? If not so, could you post more complete example?

---

<div class="post-metadata">

**Author:** ![CUDA\_VST\_RT](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@CUDA\_VST\_RT](https://forums.developer.nvidia.com/u/CUDA_VST_RT)\
**Post date:** [November 20, 2009, 9:38am UTC](https://forums.developer.nvidia.com/t/pointer-to-cudamem-using-allocated-memory-in-different-kernels/13354/3 "2009-11-20T09:38:03Z")

</div>

Ok, thanks for the fast reply.  
i’ill try to give you a overview:

c-code header:  
extern “C” void X(…, float\* a\_d);  
extern “C” void X2(…, float\*\* a\_d);

c-code:  
X(…, &a\_d);  
“sync()”  
X2(…, a\_d);

cuda.cu code:

extern “C” void X2(…, float\* a\_d)  
{  
// execute the kernel  
dim3 block(8, 8, 1);  
dim3 grid(mesh\_width / block.x, mesh\_height / block.y, 1);  
X2\_kernel\<\<\< grid, block\>\>\>(…, a\_d);

}

extern “C” void X(…, float \*\*a\_d)  
{

```
size_t MEMsize = width*height*sizeof(float);
unsigned int ArraySize = width * height;

// allocate array on Device
(cudaMalloc((void **) a_d, MEMsize));

// copy data from host to device
( cudaMemcpy(*a_d, inputBuffer, MEMsize, cudaMemcpyHostToDevice));

  //execute kernel
 X_kernel <<< nBlocks, BLOCK_DIM>>> (....., *a_d, (ArraySize));

[b](cudaMemcpy(outputBuffer, *a_d, MEMsize, cudaMemcpyDeviceToHost));

```

[/b]  
}

cuda\_kernel.cu:

**global** void X\_Kernel(…, float _a, int const size)  
{  
int Xidx = blockIdx.x_ blockDim.x + threadIdx.x;

```
if (Xidx < size)
{
	a[Xidx] = a[Xidx] * 0.5f; //for example
}

```

}

**global** void X2\_kernel(…, float\* a\_d)  
{  
unsigned int x = blockIdx.x_blockDim.x + threadIdx.x;  
unsigned int y = blockIdx.y_blockDim.y + threadIdx.y;

```
 float test = a_d[0]; //ACCESING MEM in GPU, allocatet in X()

```

…  
}

In emulation Mode i can set brakepoints an watch the values of the mem pointers in X() and X2(), and see they are the same.  
but its not working in cuda mode.

i copy some audio data and i think its working ok, because i got the correct sound output, from the first kernel X().  
but the second one isnt working alright. the adiodata copied is 2(channel) \* 1024 (blocklength) \* sizeof(float)  
that shouldnt be to much.

---

<div class="post-metadata">

**Author:** ![sagrailo](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@sagrailo](https://forums.developer.nvidia.com/u/sagrailo)\
**Post date:** [November 20, 2009, 10:37am UTC](https://forums.developer.nvidia.com/t/pointer-to-cudamem-using-allocated-memory-in-different-kernels/13354/4 "2009-11-20T10:37:50Z")

</div>

Seems to me again that everything is looking OK with regards to “a\_d” pointer handling. I’d still check for error return from all cuda\*() calls, and if there are no problems there, then it’s probably something within your second kernel code. So to debug it, it would be still good to have a minimal, but fully working code sample.

---

<div class="post-metadata">

**Author:** ![CUDA\_VST\_RT](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@CUDA\_VST\_RT](https://forums.developer.nvidia.com/u/CUDA_VST_RT)\
**Post date:** [November 20, 2009, 11:40am UTC](https://forums.developer.nvidia.com/t/pointer-to-cudamem-using-allocated-memory-in-different-kernels/13354/5 "2009-11-20T11:40:38Z")

</div>

mh the first kernel is called by a VST\_Plugin and the second one as a function of a OpenGL, which is initialized by the Plugin.

i cant give you a working example. :(

at the moment the hole system is completly frozen when executing in non emulation mode and i need to reset.

its all about this line: float test = a\_d[0]; //ACCESING MEM in GPU, allocatet in X()  
its working when i dont try to access this.

the rest of the kernel is ok.

so debugging in “cuda-mode” is a problem, because i dont get to the \_kernel.cu functions, and it cant get the values of the pointer a\_d.

thank you so much for helping me out here.

---

<div class="post-metadata">

**Author:** ![sagrailo](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@sagrailo](https://forums.developer.nvidia.com/u/sagrailo)\
**Post date:** [November 20, 2009, 11:53am UTC](https://forums.developer.nvidia.com/t/pointer-to-cudamem-using-allocated-memory-in-different-kernels/13354/6 "2009-11-20T11:53:35Z")

</div>

I guess you may already have tried this, but - can you add, after doing the sync after the first kernel completed, some host code to copy the contents of the device memory back to host, and to check is everything OK with these values (and then you could try with another cudaMalloc(), and then copy these values back to the GPU memory, but now to this freshly allocated segment of memory, so that you could somewhat judge, when the second kernel run, is it about corrupted memory, or this kernel code)? Also, if you’re on Linux, or have access to a Linux machine, can you try with CUDA debugger?

---

<div class="post-metadata">

**Author:** ![CUDA\_VST\_RT](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@CUDA\_VST\_RT](https://forums.developer.nvidia.com/u/CUDA_VST_RT)\
**Post date:** [November 20, 2009, 11:56am UTC](https://forums.developer.nvidia.com/t/pointer-to-cudamem-using-allocated-memory-in-different-kernels/13354/7 "2009-11-20T11:56:26Z")

</div>

i forgot something.

where shloud i call cudafree?

i tried to put it at the end of X2() in cuda.cu: cudafree(a\_d)

but i get a bunch of exceptions.

First-chance exception at 0x7c812afb in mulch.exe: Microsoft C++ exception: cudaError at memory location 0x0773fd2c…

---

<div class="post-metadata">

**Author:** ![CUDA\_VST\_RT](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@CUDA\_VST\_RT](https://forums.developer.nvidia.com/u/CUDA_VST_RT)\
**Post date:** [November 20, 2009, 12:02pm UTC](https://forums.developer.nvidia.com/t/pointer-to-cudamem-using-allocated-memory-in-different-kernels/13354/8 "2009-11-20T12:02:22Z")

</div>

i already copy back to the host, thats why i guess the first kernel is ok, because thats my audio stream. please look in my second post, i forgot one line in cuda.cu x() where i copy back to host.  
i will try to reallocate it in the second kernel.

i dont have access to a linux machine.

---

<div class="post-metadata">

**Author:** ![sagrailo](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@sagrailo](https://forums.developer.nvidia.com/u/sagrailo)\
**Post date:** [November 20, 2009, 12:28pm UTC](https://forums.developer.nvidia.com/t/pointer-to-cudamem-using-allocated-memory-in-different-kernels/13354/9 "2009-11-20T12:28:19Z")

</div>

I’d say calling cudaFree() after X2() call completed (or at the end of X2(), albeit the whole practice of allocating memory in one routine, and freeing it some other routine isn’t particularly nice) is the right way to go (but do check that you don’t have multiple cudaFree() calls around). Further - I don’t have much Windows experience, but seems like it is indeed something about memory corruption; if you skip your second kernel call, and cudaFree() is working OK in that case, then at least you’ll know that second function/kernel is messing something up; also, if you could get to the cudaFree() error number in the case of that exception throw, then it would be something to start with…

---

<div class="post-metadata">

**Author:** ![CUDA\_VST\_RT](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@CUDA\_VST\_RT](https://forums.developer.nvidia.com/u/CUDA_VST_RT)\
**Post date:** [November 23, 2009, 9:01am UTC](https://forums.developer.nvidia.com/t/pointer-to-cudamem-using-allocated-memory-in-different-kernels/13354/10 "2009-11-23T09:01:40Z")

</div>

mh i dont know, now i am making a new cudamemcpy in the second kernel, thats working.  
using the pointer to an existing context, doesnt work for me :(  
(though it does work in emulation mode)

it seems that the pointer a\_d gets corrupted.