# Not working correctly new () and malloc () inside the kernel, why?

**URL:** <https://forums.developer.nvidia.com/t/not-working-correctly-new-and-malloc-inside-the-kernel-why/32757>\
**Category:** CUDA Programming and Performance\
**Created:** [April 2, 2014, 1:37pm UTC](https://forums.developer.nvidia.com/t/not-working-correctly-new-and-malloc-inside-the-kernel-why/32757 "2014-04-02T13:37:00Z")\
**Posts on this page:** 3\
**Page:** 1

<div class="post-metadata">

**Author:** ![juravlev\_work](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@juravlev\_work](https://forums.developer.nvidia.com/u/juravlev_work)\
**Post date:** [April 2, 2014, 1:37pm UTC](https://forums.developer.nvidia.com/t/not-working-correctly-new-and-malloc-inside-the-kernel-why/32757/1 "2014-04-02T13:37:00Z")

</div>

**global** void gpu\_test\_malloc(int \* a\_t\_size)  
{

```
	int64 t_size=a_t_size[0];//in 40435848;
	void * t_void=NULL;
	label_test:		
	t_void=malloc(t_size);
	if(t_void==NULL)
	{
		t_size--;
		goto label_test;
	};
	a_t_size[1]=t_size;//out 7674032

```

};

…

```
//test
int * host_mem=new int[2];
int * device_mem;
host_mem[0]=40435848;//need
cudaStatus=cudaMalloc(&device_mem,sizeof(int)*2);
if (cudaStatus != cudaSuccess)throw;

cudaStatus=cudaMemcpy(device_mem,host_mem,sizeof(int)*2,cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess)throw;

gpu_test_malloc<<<1,1>>>(device_mem);

cudaStatus=cudaDeviceSynchronize();
if (cudaStatus != cudaSuccess)throw;

cudaStatus=cudaMemcpy(host_mem,device_mem,sizeof(int)*2,cudaMemcpyDeviceToHost);
if (cudaStatus != cudaSuccess)throw;

```

…  
//===============  
host\_mem[0]=40435848  
host\_mem[1]= 7674032  
//===============  
GTX Titan(driver 320.14)  
Nsight 3.0.0.13123  
CUDA 5.0  
64bit,compute\_30,sm\_30  
//==============

---

<div class="post-metadata">

**Author:** ![juravlev\_work](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@juravlev\_work](https://forums.developer.nvidia.com/u/juravlev_work)\
**Post date:** [April 3, 2014, 8:56am UTC](https://forums.developer.nvidia.com/t/not-working-correctly-new-and-malloc-inside-the-kernel-why/32757/2 "2014-04-03T08:56:31Z")

</div>

#include “cuda\_runtime.h”  
#include “device\_launch\_parameters.h”

#include \<stdio.h\>

**global** void gpu\_test\_malloc(int \* a\_t\_size)  
{

```
__int64 t_size=a_t_size[0];//40435848;
void * t_void=NULL;
label_test:
t_void=malloc(t_size);
if(t_void==NULL)
{
	t_size--;
	goto label_test;
};
a_t_size[1]=t_size;

```

};

int main()  
{

cudaError\_t cudaStatus;  
try  
{  
int \* host\_mem=new int[2];  
int \* device\_mem;  
host\_mem[0]=40435848;//это нужный мне объём  
cudaStatus=cudaMalloc(&device\_mem,sizeof(int)\*2);  
if (cudaStatus != cudaSuccess)throw;  
cudaStatus=cudaMemcpy(device\_mem,host\_mem,sizeof(int)\*2,cudaMemcpyHostToDevice);  
if (cudaStatus != cudaSuccess)throw;

```
	gpu_test_malloc<<<1,1>>>(device_mem);

	cudaStatus=cudaDeviceSynchronize();
	if (cudaStatus != cudaSuccess)throw;
	cudaStatus=cudaMemcpy(host_mem,device_mem,sizeof(int)*2,cudaMemcpyDeviceToHost);
	if (cudaStatus != cudaSuccess)throw;
	cudaStatus=cudaFree(device_mem);
	return cudaStatus;
}catch(...)
{
	return cudaStatus;
};

```

# }

# GTX Titan(драва 335.23) Nsight 3.2.2.13351 CUDA 5.5 64bit,compute\_35,sm\_35

# =============== host\_mem[0]=40435848 host\_mem[1]= 7674032

---

<div class="post-metadata">

**Author:** ![vyas](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@vyas](https://forums.developer.nvidia.com/u/vyas)\
**Post date:** [April 4, 2014, 5:49pm UTC](https://forums.developer.nvidia.com/t/not-working-correctly-new-and-malloc-inside-the-kernel-why/32757/3 "2014-04-04T17:49:53Z")

</div>

The device side malloc() is a suballocator from a fixed size heap. Given that the kernel loop is continuously allocating memory, I suspect the backing heap is being exhausted. The size of the heap can be changed using the CUDA API :

```auto
cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size)

```

. More information about device side malloc() can be obtained in the toolkit documentation : [http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#dynamic-global-memory-allocation-and-operations](http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#dynamic-global-memory-allocation-and-operations)