Code with thrust::remove_if fails on CUDA 7.5 but runs well in CUDA 6.5

[metorm]

The following code fails and throws thrust::system_error under CUDA 7.5 & VS 2013 community & win10 & GT750M

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <thrust/system_error.h>

#include <thrust/device_vector.h>
#include <thrust/functional.h>
#include <thrust/extrema.h>

#include <thrust/remove.h>
#include <thrust/copy.h>
#include <thrust/fill.h>
#include <thrust/sequence.h>

#include <thrust/execution_policy.h>

#include <iostream>

int main()
{
    thrust::device_vector<int> AA(10, 1);
    thrust::sequence(AA.begin(), AA.end());
    thrust::device_vector<bool> SS(10, false);
    thrust::fill(SS.begin(), SS.begin() + 5, true);
    try{
        thrust::device_vector<int>::iterator new_end = thrust::remove_if(thrust::device,
            AA.begin(), AA.begin()+3, SS.begin(), thrust::identity<bool>());
    }
    catch (thrust::system_error &e)
    {
        std::cerr << e.code() << e.what() << std::endl;
        exit(-1);
    }

    return 0;
}

The exception detail was cuda:11  invalid argument. But the code runs well on another computer which runs CUDA 6.5.

[metorm]

Well, when I compiled the above code as x32 program under Win10 x64, the program throws thrust::system_error , but if I compile it as x64 program, everything was OK.
I don't know if thrust is designed like this or it is a bug.

[stic]

Could it be linked to the fact that trivial_copy.inl calls trivial_copy_detail::checked_cudaMemcpyAsync with cudaMemcpyDefault which might not work within 32-bit process that doesn't support unified addressing?
http://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__UNIFIED.html#axzz4GAddb5Rf

[metorm]

@stic I think so. Thanks for providing this detailed information. Before this I only know newer versions of CUDA and thrust sometimes fails in 32-bit programs.

[jaredhoberock]

@egaburov: Can you say whether this is an effect of limited 32b support on Windows, and if so, is it likely we can fix this?

[3gx]

This is fixed in CUDA 8.0; it doesn't repro with cuda-next-release and 1.8.3 branches.

[jaredhoberock]

@egaburov: Thanks! Does the problem exist in the master branch? Just trying to figure out whether this issue needs to be closed.

[3gx]

I do not observe failure in the master branch either. The issue can be closed.