selectively compile headers and class functions in CUDA

Question

I am attempting to utilize my c++ classes within CUDA.

I have a class as such:

#include<string>
#include<stdlib.h>

class exampleClass{
int i;
__host__ __device__ exampleClass(int _i):i(_i){};
__host__ __device__ void increment(){i++;}
__host__ __device__ string outputMessage(return itoa(i);}

};

I have set this in a .cu file and set to compile CUDA c/c++

This fails to compile with nvcc because cuda doesn't have strings.

What I'd like to do is retain the CUDA only functions by doing something like:

#ifndef __CUDA_ARCH__
  #include<string>
#endif
    #include<stdlib.h>

    class exampleClass{
    int i;
    __host__ __device__ exampleClass(int _i):i(_i){};
    __host__ __device__ void increment(){i++;}
#ifndef __CUDA_ARCH__
     string outputMessage(return itoa(i);}
#endif

    };

But I know this doesn't work...at least, it isn't working for me. The nvcc doesn't like the string inclusion nor, obviously, the function that requires the string type.

Apologies if the example isn't top-notch. In summary, what I'd like to do is have core class members executable on CUDA while maintaining the ability to have fancy host operations for analysis and output on the host side.

UPDATE: My end goal here is to have a base class, containing several pointer types to several polymorphic classes. This base class itself is going to be derivable. I thought this was possible in CUDA5.0. Am I mistaken?

I've had no problem using stl classes, such as string or vector, with nvcc. You certainly can't send the string to a kernel, or use the string function in a kernel, but that doesn't look like that's what you are trying to do...you should just be able to remove the "__device__" decoration. — IdeaHat, Mar 28 '13 at 20:03
MadScienceDreams was correct in the answer to my question. The problem underlying the problem, however, found its solution here:http://stackoverflow.com/questions/5994005/cuda-external-calls-not-supported — PaulD, Mar 28 '13 at 22:50
More particularly, you need to include the -dc flag in the CUDA C/C++ compiler to allow linkage between multiple cu files. — PaulD, Mar 28 '13 at 22:53

score 2 · Accepted Answer · answered Mar 28 '13 at 20:16

2

The following code builds, though I didn't run it:

class exampleClass{
int i;
public:
__host__ __device__ exampleClass(int _i):i(_i){};
__host__ __device__ void increment(){i++;}

 __host__ string outputMessage(){ return "asdf";}


};

__global__ void testkernel (                        
    exampleClass *a,
    int IH, int IW)
{
    const int i = IMUL(blockIdx.x, blockDim.x) + threadIdx.x;
    const int j = IMUL(blockIdx.y, blockDim.y) + threadIdx.y;


    if (i<IW && j<IH) 
    {
        const int i_idx = i + IMUL(j, IW);  
        exampleClass* ptr = a+i_idx;
        ptr->increment();
    }
}

__host__ void test_function(exampleClass *a,
    int IH, int IW)
{
    for (int i = 0; i < IW; i++)
        for (int j = 0; j < IH; j++)
        {
            const int i_idx = i + j*IW;
            exampleClass* ptr = a+i_idx;
            cout << ptr->outputMessage();
        }
}

Note that you'll have to move the classes from device to host memory for this to "work" properly. If you try to do anything fancy with the classes (such as polymorphism, for example) this will probably blow up.

answered Mar 28 '13 at 20:16

IdeaHat

7,641
1
22
53

I thought polymorphism was allowed in CUDA5.0? My "exampleClass" is a base class with several derivatives. As well, the exampleClass contains several points that hold polymorphic classes. – PaulD Mar 28 '13 at 21:07
@PaulD. Huh, that's a cool change from 4.0. CUDA is one of those few software packages that seems to add nice features every year! The VTable of these classes must have a reference to both the `__host__` and `__device__` function pointers, so the class size will grow pretty quickly. Note that virtual inheritance is still unsupported. – IdeaHat Mar 29 '13 at 12:16
I'm unfamiliar with VTable...I'm more than a little self taught. Can you explain that a bit or post a link? As well, the nVidia website says that virtual inheritance has been supported since Toolkit 4.0. https://developer.nvidia.com/cuda-toolkit-40 – PaulD Mar 29 '13 at 15:22
The code standard 5 says ["It is not allowed to pass as an argument to a __global__ function an object of a class derived from virtual base classes."](http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#restrictions), so no virtual inheritance. It also states "The virtual function table is placed in global or constant memory by the compiler." though I have no idea how that happens (static load?). [VTable in wikipedia](en.wikipedia.org/wiki/Virtual_method_table) – IdeaHat Mar 29 '13 at 15:32
I was under the impression that a "virtual base class" was a particular type of inheritance: http://stackoverflow.com/questions/21558/in-c-what-is-a-virtual-base-class – PaulD Apr 01 '13 at 13:45
Yeah, it is "virtual inheritance". http://en.wikipedia.org/wiki/Virtual_inheritance – IdeaHat Apr 01 '13 at 13:55

selectively compile headers and class functions in CUDA

1 Answers1