tagged_iterator does not propagate its tag to its nested pointer & reference types

consider

```
int *raw_ptr;
auto ptr = reinterpret_tag(raw_ptr);
```

The result of `*ptr` is `int &` when it should be `reference`.

The result of `&*ptr` is `int *` when it should be `pointer`.

After this is fixed, clean up the implementation of `cuda::detail::fill_n` so that it doesn't have to use `thrust::pointer` directly.

Forwarded from http://code.google.com/p/thrust/issues/detail?id=462