CUDA Memory Access Violation when creating an object in kernel function

Hi everyone,

I’m developing some graphs algorithms on cuda, I have to work with large amounts of data, so to allocate a large memory in CUDA kernel function using operator new I set the value of cudaLimitMallocHeapSize to the size of free device memory (~1.7-1.8G at every launch). But if I try to create an object in kernel function I get a memory access violation in the line

i = 42;

throwed by memory checker. Here is the code:

class Test
{
private:
    int i;

public:
    __device__ Test()
    {
        i = 42;
    }
};

__global__ void test()
{
    Test *m = new Test();
    if (m == NULL)
    {
        printf("m == NULL\n");
    }
    else
    {
        printf("OK\n");
    }
}

int main(int argc, char *argv[])
{
    size_t free;
    size_t total;
    cudaMemGetInfo(&free, &total);
    cudaCheckError(cudaDeviceSetLimit(cudaLimitMallocHeapSize, free));

    test << <1, 1 >> >();
    cudaCheckError(cudaDeviceSynchronize());
    cudaCheckError(cudaDeviceReset());
    return 0;
}

But after this exception object

Test *m

is a normal object with address in device memory and with its field i is equal to 42. If I remove the constructor of Test class, the same exception raised in the cuda source file device_functions.h in the function

static __forceinline__ void* memset(void *dest, int c, size_t n)
{
  __nvvm_memset((unsigned char *)dest, (unsigned char)c, n, /*alignment=*/1);
  return dest;
}

If I set the heap size to 110241024*1024 = 1G - the same exception occures.
750 MB - no exception.
900 MB - no exception.
1000 MB - no exception.
1020 MB - exception.

Why it happens? Is there any limit to heap size lower then the size of available memory?

I’m using Visual Studio Ultimate 2013 with Nsight 4.1.
The device is GeForce GTX 650 Ti with compute capability 2.0.

Thanks for any help.

UPD: 2GB DDR5 memory total on device. OS is Windows 8.1 Professional x64.

__
Best Regards,
Alexander Golovkov

Hi Alexander,

We have reproduced your problem on our side. Developers are investigating it. Thank you for patience and sorry the the inconvenience.

Hi qzhang,

Thanks for your reply.
Please let me know if there will be any solutions.

__
Best Regards,
Alexander Golovkov