0

Given the following code:

#define DEVICE __device__
#define HOST __host__

#include <thrust/host_vector.h>
#include <thrust/device_vector.h>

template <typename T, std::size_t N>
class Container {
public:

    DEVICE HOST
    Container() {

    }

private:
    T data[N];
};

typedef Container<double, 7> double7;

template <std::size_t N = 10 >
class History {
public:

    DEVICE HOST
    History() {

    }

    DEVICE HOST
    virtual ~History() {

    }

private:

    double7 history[N];
};

int main() {

    try {

        thrust::host_vector<History<> > histories(1);
        thrust::device_vector<History<> > d_histories = histories;
    } catch (const thrust::system_error &) {
        std::cerr << "boo boo" << std::endl;
    }

    return 0;
}

The above runs without error on toolkit release nvcc 5.0, v 0.2.1221 and fails on nvcc 5.5 v5.5.0. I realize virtual functions in __global__ are not supported as the documentation clearly states.

My question is this: what changed between the two versions to "break" virtual functions? Was this an accidental omission in 5.0 that was corrected in 5.5?

Robert Crovella
  • 143,785
  • 11
  • 213
  • 257
Tyler Jandreau
  • 4,245
  • 1
  • 22
  • 47
  • 1
    The difference appears to be in the change from thrust v1.5.3 in CUDA 5 to thrust v1.7 in CUDA 5.5. If I run your code as-is in CUDA 5 with thrust v1.5.3 (the version of thrust that ships with CUDA 5.5) it runs without the system_error. If I upgrade thrust to v1.7 but otherwise use CUDA 5, then it results in the system_error. The system under the earlier version of thrust may have used a different method to copy objects. In any event, I don't think the CUDA limitation was any different between CUDA 5 and CUDA 5.5 in this respect, but the behavior of thrust was different. – Robert Crovella Sep 26 '14 at 03:36
  • Sorry I meant to say that thrust v1.5.3 is the version of thrust that ships with CUDA **5** not **5.5** – Robert Crovella Sep 26 '14 at 12:37
  • @RobertCrovella Thank you for your answer. I knew I was not crazy when my code ran on one machine and not the other! Newer versions of thrust must be more stringent with the `virtual` keyword. That's a good lessons-learned for future development. – Tyler Jandreau Sep 26 '14 at 12:45

0 Answers0