Allocating Pointer to Pointer inside a Kernel?

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello,

Here is the following code where I am stuck at (-lineinfo) and I want to know if the double pointer allocation is allowed within a cuda kernel.

__global__ void LAUNCHBOUNDS(MY_KERNEL_MAX_THREADS) Kernel()
{
	double** matrix = (double**)malloc(6 * sizeof(double*)),
		** invmatrix = (double**)malloc(6 * sizeof(double*));
	for (int i = 0; i < 6; i++)
	{
		matrix[i] = (double*)malloc(6 * sizeof(double));
		invmatrix[i] = (double*)malloc(6 * sizeof(double));
	}
	matrix[0][0] = 1.0; matrix[0][1] = 1.0; matrix[0][2] = 1.0; matrix[0][3] = 1.0; matrix[0][4] = 1.0; matrix[0][5] = 1.0;
	matrix[1][0] = 1.0; matrix[1][1] = 1.0; matrix[1][2] = 1.0; matrix[1][3] = 1.0; matrix[1][4] = 1.0; matrix[1][5] = 1.0;
	matrix[2][0] = 1.0; matrix[2][1] = 1.0; matrix[2][2] = 1.0; matrix[2][3] = 1.0; matrix[2][4] = 1.0; matrix[2][5] = 1.0;
	matrix[3][0] = 1.0; matrix[3][1] = 1.0; matrix[3][2] = 1.0; matrix[3][3] = 1.0; matrix[3][4] = 1.0; matrix[3][5] = 1.0;
	matrix[4][0] = 1.0; matrix[4][1] = 1.0; matrix[4][2] = 1.0; matrix[4][3] = 1.0; matrix[4][4] = 1.0; matrix[4][5] = 1.0;
	matrix[5][0] = 1.0; matrix[5][1] = 1.0; matrix[5][2] = 1.0; matrix[5][3] = 1.0; matrix[5][4] = 1.0; matrix[5][5] = 1.0;
	for (int i = 0; i < 6; i++)
	{
		free(matrix[i]);
		free(invmatrix[i]);
	}
	free(matrix);
	free(invmatrix);
}

Thanks

2

Yes, double pointer allocation is allowed in a kernel

3

Ok. But why i got a -lineinfo pointed to matrix[i] = (double*)malloc(6 * sizeof(double));
. It seems to be a memory issue i got from there.

4

I tried to use compute-sanitizer --tool memcheck AppName and All i got as errors is : Internal Sanitizer Error: The Sanitizer failed to handle a hardware exception.

5

Here is what i got from Visual Studio :
image

6

You need to check if malloc returns nullptr.
The device heap has limited size. Its size can be set via cudaDeviceSetLimit (CUDA Runtime API :: CUDA Toolkit Documentation)

If you need more help, please post a minimal complete reproducer for your problem.

7

All right, i will try it later. But will the increasing limit impact the Kernel Performance?