CUDA's get_value and assign_value WAR doesn't work for msvc 2005

[GoogleCodeExporter]

MSVC 2005 produces multiply-defined symbols due to this workaround.

The WAR for the WAR is to put the implementation of get_value inside an 
anonymous namespace to force its visibility to the current translation unit. 
Then the get_value facade calls the version in the anon namespace.

Original issue reported on code.google.com by jaredhoberock on 22 Feb 2012 at 1:45

[GoogleCodeExporter]

Original comment by jaredhoberock on 22 Feb 2012 at 8:39

• Changed title: CUDA's get_value and assign_value WAR doesn't work for msvc 2005

[GoogleCodeExporter]

using an anon namespace makes the unit tester go into lala land. so we'll have 
to try a slightly different WAR

Original comment by jaredhoberock on 22 Feb 2012 at 9:50

[GoogleCodeExporter]

Didn't work.  Seems that nvcc doesn't correctly register the kernels launched 
by get_value & assign_value, and an exception is thrown when the kernel gets 
launched.

I don't know how to WAR this.  Seems to be some broken interaction between nvcc 
& msvc 2005.  Let's unsupport msvc 2005 in 1.6.  We can revisit a WAR in 1.7 if 
there's interest.

Original comment by jaredhoberock on 23 Feb 2012 at 12:44

• Added labels: Milestone-Release1.7
• Removed labels: Milestone-Release1.6

[GoogleCodeExporter]

I'm going to commit a change which allows code to build without a compiler 
error (due to msvc 2005 not evaluating enable_if_convertible correctly).

The code won't necessarily link together correctly, and even if it does, it may 
throw an exception at runtime: "invalid device function" because cudafe doesn't 
correctly register all kernels with the runtime.

I suggest in 1.6 that we add a warning if _MSC_VER <= 1400 (msvc 2005), but not 
be so heavy-handed as to issue an #error.

Original comment by jaredhoberock on 24 Feb 2012 at 1:24

[GoogleCodeExporter]

The root of the problems seem to stem from MSVC 2005 not correctly evaluating 
thrust::detail::pointer_raw_pointer for wrapped pointers inside Thrust's 
implementation. This metafunction is correctly evaluated in code external to 
Thrust which #includes the proper headers. I suspect at some points of 
instantiation, the type of wrapped pointer may not be complete yet.

I also noticed another issue with enable_if_pointer_is_convertible, but my 
suspicion is that the root cause is the same, as 
enable_if_pointer_is_convertible relies on pointer_raw_pointer.

It looks like we can WAR the problem by providing specializations of 
pointer_raw_pointer for each type of problematic pointer wrapper. If so, we can 
support MSVC 2005 in 1.6.

Original comment by jaredhoberock on 29 Feb 2012 at 1:55

[GoogleCodeExporter]

Original comment by jaredhoberock on 29 Feb 2012 at 11:09

• Added labels: Milestone-Release1.6
• Removed labels: Milestone-Release1.7

[GoogleCodeExporter]

This issue was closed by revision d36b48b87ea6.

Original comment by jaredhoberock on 1 Mar 2012 at 1:41

• Changed state: Fixed