__host__ and __device__ qualifies

Accelerated Computing CUDA CUDA Programming and Performance
1

He everyone, I’m having trouble converting a ant algorithm to cuda and was wondering if anyone could help me. We have an existing ant algorithm written in c++ used for finding solutions to the degree constrained minimum spanning tree problem. I’m trying to gradually convert the algorithm over to cuda in chunks so I altered all necessary classes to be structs with associated function calls that take a struct* as their first parameter. Pretty straightforward.

Now I need to have those associated struct apis callable on either the host or the device side. In the documentation it says that I can mix device and host specifiers but I can’t get the first one to compile. I’m always getting an error

dcmst_kernels.cu(14): error: calling a host function from a __device__/__global__ function is only allowed in device emulation mode

Here are the relevant sections of my code:

This is the only function I’ve tried to call on the device side so far. It’s in the file graph.cu which was renamed from graph.cpp. It’s prototyped in graph.cuh like so: “int Graph_getNumVertexes(Graph *graph)” I thought maybe I have to add the device specifier there but then the c++ code included the file complains about device not naming a type.

__device__ __host__

int Graph_getNumVertexes(Graph *graph)

{

	return graph->numVertexes;

}

This is the kernel that I’m trying to call the function from.

#include "graph.cuh"

#include "cuda_macros.h"

#include "mt19937_ref.cu"

__global__

void AntsMove_Kernel(Graph* graph, int steps)

{

   int numVertexes = Graph_getNumVertexes(graph);

   if(LINEAR_THREAD_INDEX_GRD() < numVertexes) {

}

   uint rand = mt19937gl();

}

The relevant portions of my makefile and the output from Make when I try to compile it.

graph.o: graph.cuh graph.cu

	$(CUDA_CC) $(CXXFLAGS) -c graph.cu

dcmst_kernels.o : dcmst_kernels.cu cuda_macros.h mt19937_ref.cu

	$(CUDA_CC) $(CXXFLAGS) -c dcmst_kernels.cu

nvcc -g -lcudart -L/usr/local/cuda/lib/ -L/raid1/thesis_project/NVIDIA_GPU_Computing_SDK/C/lib/ -c graph.cu

nvcc -g -lcudart -L/usr/local/cuda/lib/ -L/raid1/thesis_project/NVIDIA_GPU_Computing_SDK/C/lib/ -c dcmst_kernels.cu

dcmst_kernels.cu(14): error: calling a host function from a __device__/__global__ function is only allowed in device emulation mode

1 error detected in the compilation of "/tmp/tmpxft_0000281a_00000000-4_dcmst_kernels.cpp1.ii".

Any help or ideas would be appreciated. I assume it’s something I don’t know about the build process. Thanks.

2

Not every mathematical function is supported on the device. You may find appendix B-1 of the CUDA programming guide 2.0 most enlightening…