Thrust exception: "thrust::system::system_error at memory location 0x00000000"

Question

I wrote this code of a CUDA Kernel assign() using the class device_vector for initializing a vector. This kernel is launched by a class member function as a solution to the question:

CUDA kernel as member function of a class

and according to

https://devtalk.nvidia.com/default/topic/573289/mixing-c-and-cuda/.

I'm using a GTX650Ti GPU, Windows 8.1, Visual Studio 2013 Community and CUDA Toolkit 7.5.

The code initTest.cu does compile but an exception is thrown making reference to the file trivial_copy.inl.

"First-chance exception at 0x775B5B68 in initTest.exe: Microsoft C++ exception: thrust::system::system_error at memory location 0x0116F3C8. If there is a handler for this exception, the program may be safely continued."

Does anyone know why this problem occurs?

The header file foo.cuh is:

#ifndef FOO_CUH
#define FOO_CUH
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <thrust/device_vector.h>
#include <vector>
using namespace thrust;
using namespace std;

__global__ void assign(float *x, const float &constant, const unsigned int &n)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < n)
        x[i] = constant;
}
class foo
{
    public:
    foo(const unsigned int &);
    void init(const float &);
    vector<float> domain;
private:
    unsigned int samples;
};
foo::foo(const unsigned int &n)
{
    vector<float> result(n);
    domain = result;
    samples = n;
}
void foo::init(const float &value)
{
    device_vector<float> result(samples);
    assign <<< 1, domain.size() >>>(raw_pointer_cast(result.data()), value, samples);
    thrust::copy(result.begin(), result.end(), domain.begin());
}
#endif

and the main function defined in initTest.cu is:

#include "foo.cuh"
#include <iostream>

int main()
{
    foo a(10);
    a.init(0.5);
    for (unsigned int i = 0; i < a.domain.size(); i++)
    {
        if (i == 0)
            cout << "{ ";
        else if (i == a.domain.size() - 1)
            cout << a.domain[i] << " }";
        else
            cout << a.domain[i] << ", ";
    }
    cin.get();
    return 0;
}

Answer 1

This is illegal:

__global__ void assign(float *x, const float &constant, const unsigned int &n)
                                             ^                             ^

Kernel parameters cannot be pass-by-reference.

When I remove the ampersands:

__global__ void assign(float *x, const float constant, const unsigned int n)

Your code runs correctly for me.

I would suggest you use proper cuda error checking. Doing so would have focused your attention on the kernel. Instead, the error was uncaught until thrust detected it and threw a system_error, which doesn't help to identify the source of the error.