4

I'm considering the following simple code in which I'm converting thrust::host_vector<int>::iterator h_temp_iterator = h_temp.begin(); and thrust::device_vector<int>::iterator d_temp_iterator = d_temp.begin(); to raw pointers.

To this end, I'm passing &(h_temp_iterator[0]) and &(d_temp_iterator[0]) to a function and a kernel, respectively. The former (CPU case) compiles, the latter (GPU case) not. The two cases should be in principle symmetric, so I do not understand the reason for the error message which is:

Error   1   error : no suitable conversion function from "thrust::device_ptr<int>" to "int *" exists

The configurations are:

  1. Windows 7, Visual Studio 2010, CUDA 7.5, compiling for the 3.5 architecture.
  2. Windows 10, Visual Studio 2013, CUDA 8.0, compiling for the 5.2 architecture.

CODE

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

__global__ void testKernel(int *a, const int N)
{
    int i = threadIdx.x;

    if (i >= N) return;

    a[i] = 2;
}

void testFunction(int *a, const int N)
{
    for (int i = 0; i < N; i++) a[i] = 2;
}

int main()
{
    const int N = 10;

    thrust::host_vector<int> h_temp(N);
    thrust::device_vector<int> d_temp(N);

    thrust::host_vector<int>::iterator h_temp_iterator = h_temp.begin();
    thrust::device_vector<int>::iterator d_temp_iterator = d_temp.begin();

    testFunction(&(h_temp_iterator[0]), N);
    testKernel<<<1, N>>>(&(d_temp_iterator[0]), N);

    for (int i = 0; i < N; i++) printf("%i %i\n", i, h_temp[i]);

    return 0;
}
Vitality
  • 20,705
  • 4
  • 108
  • 146
  • 1
    This is just a very mild variation on this, isn't it ? https://stackoverflow.com/q/11113485/681865 – talonmies Jul 27 '17 at 13:01
  • 2
    In the thrust tag based typed model (the original model), any device memory is represented by a `thrust::device_ptr`. This has pointer semantics, but isn't a pointer, which helps enforce host/device pointer type safety. In your code above, the iterator decays to a `device_ptr` not a pointer. You must cast the `device_ptr` to get a raw pointer you can pass to device code – talonmies Jul 27 '17 at 13:28
  • Thank you, talonmies, for your tireless help. I have solved the problem. Would you like to post an answer? – Vitality Jul 27 '17 at 13:44
  • You could go ahead and write your own if you like. I'll upvote it. – talonmies Jul 27 '17 at 13:46
  • Thanks. I have uploaded an answer with a fully working code. – Vitality Jul 27 '17 at 14:45
  • This is an example of a question and answer which are better served by no MCVE. It just takes up screen space IMHO and the issue is clear with merely the kernel invocation one-liner. – einpoklum Jul 27 '17 at 14:54

2 Answers2

5

Following talonmies' comments, the solution is to pass

thrust::raw_pointer_cast(&d_temp_iterator[0])

and not

&d_temp_iterator[0]

In the following, the fully working code

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

__global__ void testKernel(int *a, const int N)
{
    int i = threadIdx.x;

    if (i >= N) return;

    a[i] = 2;

    printf("GPU %i %i\n", i, a[i]);
}

void testFunction(int *a, const int N)
{
    for (int i = 0; i < N; i++) {
        a[i] = 2;
        printf("CPU %i %i\n", i, a[i]);
    }
}

int main()
{
    const int N = 10;

    thrust::host_vector<int> h_temp(N);
    thrust::device_vector<int> d_temp(N);

    thrust::host_vector<int>::iterator h_temp_iterator = h_temp.begin();
    thrust::device_vector<int>::iterator d_temp_iterator = d_temp.begin();

    int *temp = thrust::raw_pointer_cast(&d_temp_iterator[0]);

    testFunction(&(h_temp_iterator[0]), N);
    testKernel<<<1, N>>>(temp, N);

    return 0;
}
Vitality
  • 20,705
  • 4
  • 108
  • 146
3

As you are using thrust, another prettier solution is to get the pointer with data() and cast it to a raw pointer:

thrust::raw_pointer_cast(d_temp_iterator.data())
luke8800gts
  • 398
  • 3
  • 7