Error: External calls are not supported (found non-inlined call to _Znwm)

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi all… I am new to CUDA. I am getting error

quadtree.cu: In function `_ZN10quadsquare12AddHeightMapERK14quadcornerdata’:
quadtree.cu:45: warning: cast to pointer from integer of different size
./quadtree.cu(45): Error: External calls are not supported (found non-inlined call to _Znwm)

when i am compiling program. My program :

#include<stdio.h>

typedef unsigned short uint16;

class quadsquare;

device struct quadcornerdata {

const quadcornerdata*	Parent;

quadsquare*	Square;

int	ChildIndex;

int	Level;

int	xorg, zorg;

};

device struct quadsquare {

quadsquare*	Child[4];

uint16	Error[6];	// e, s, children: ne, nw, sw, se

uint16	MinY, MaxY;	// Bounds for frustum culling and error testing.

unsigned char	EnabledFlags;	// bits 0-7: e, n, w, s, ne, nw, sw, se

__device__ __inline__ quadsquare(quadcornerdata* pcd);
__device__ __inline__ void AddHeightMap(const quadcornerdata& cd);

};

quadsquare *root = NULL;

quadcornerdata *RootCornerData;

device quadsquare::quadsquare(quadcornerdata* pcd)

{
pcd->Square = this;
}

device void quadsquare::AddHeightMap(const quadcornerdata& cd)
{

for (int i = 0; i < 4; i++) {

	quadcornerdata	q;
	
	if (Child[i] == NULL) {

		[b]Child[i] = new quadsquare(&q); // Error about this line.[/b]

	}	
}

}

global void testKernel(quadsquare *root,quadcornerdata *RootCornerData)
{

root->AddHeightMap(*RootCornerData);

}

int main()
{
RootCornerData=(struct quadcornerdata *)malloc(sizeof(struct quadcornerdata));

*RootCornerData = (struct quadcornerdata) {NULL,NULL,0,15,0,0};


quadsquare *p;

cudaMalloc(&p, sizeof(quadsquare));
cudaMemcpy(p,root, sizeof(quadsquare), cudaMemcpyHostToDevice);


quadcornerdata *q;


cudaMalloc(&q, sizeof(quadcornerdata));
cudaMemcpy(q,RootCornerData, sizeof(quadcornerdata), cudaMemcpyHostToDevice);


testKernel<<< 1, 1 >>>(p,q);

}

I searched in forums and tried to compile as : nvcc quadtree.cu -arch=sm_20 and nvcc quadtree.cu -arch=sm_12

But , Still getting error. Can any one please tell the error what about it is and how to solve the problem ?

2

Looks like the compiler doesn’t like having non-inlined constructors in classes.

3

Hi… Thanks for reply. I have some doubts :

  1. In program, I made constructors as inlined functions. It’s showing same error. How to solve this problem?
  2. Can we allocate memory dynamically in device function using new?

Can You please clarify these?

4

When I said inline, I was meaning declare all member functions inside the class declaration itself. I have managed to get simple classes with trivial constructors to compile like this in CUDA 3.2 and 4.0rc.

[

Yes, but only on a Fermi GPU when code is compiled for the sm_20 or sm_21 target architecture and using the CUDA 4.0rc toolkit. malloc might work in CUDA 3.2 as well, but new and delete are CUDA 4 features AFAIK.

5

Thanks for your clarification…

I used this structure decleration in the program:

class quadsquare;

device struct quadcornerdata {

const quadcornerdata* Parent;
quadsquare* Square;
int ChildIndex;
int Level;
int xorg, zorg;

};

device struct quadsquare {

quadsquare* Child[4];
uint16 Error[6]; // e, s, children: ne, nw, sw, se
uint16 MinY, MaxY; // Bounds for frustum culling and error testing.
unsigned char EnabledFlags; // bits 0-7: e, n, w, s, ne, nw, sw, se

device inline quadsquare(quadcornerdata* pcd);
device inline void AddHeightMap(const quadcornerdata& cd);

};

I declared the functions inside class decleration only. Still getting error.

Can you please show me the correct decleration ?

6

Perhaps my terminology was a bit confusing in the earlier replies. Inline the definition of all class methods inside the definition of the class.

So something like

__device__ struct someclass {

    __device__ someclass() { .. constructor code .. };

    __device__ somemethod() { .. method code .. };

};
7

Hi… Now I made those function definations inside the structure only. Then also getting same error. Here I attached my program.

Can you please tell me where i have to made changes and why i am still getting error?
quadtree.cu (1.38 KB)

8

That code compiles without error (a couple of warnings but nothing important) using the cuda 4.0rc compiler with gcc 4.3 on linux.

9

Thanks you vary much . I will try to install CUDA 4.0

10

Hi… I tried to download the CUDA 4.0 . I registered but Website not generating any password to my email. Can you please tell the procedure for download ?

11

You will just have to wait. Registration is not automatic. If your application to register as a developer is accepted, you will get a password emailed to you.