Add support for CUDA_NO_HALF

[georgelyu]

Sometimes the half type defined in CUDA library conflicts with the one defined in other libraries (in Imath for example, which is the problem I'm encountering right now). I wanna use #define CUDA_NO_HALF to get past this, but the type is used in array.hpp. So, I wrote a naive way to support the macro, which is checked in Line 2535 of cuda_fp16.hpp.

[eyalroz]

Hello Chaoyang,

First - thanks for the interest in contributing. :-)

Personally, I don't do half-precision work, so my experience here is very limited. I noticed that there are several relevant definitions in the CUDA headers:

• CUDA_NO_HALF
• __CUDA_NO_HALF_OPERATORS__
• __CUDA_NO_HALF_CONVERSIONS__

Are you sure we should be using the first of these?

Also, would it not be sufficient if I just used __nv_half instead of half?

[eyalroz]

How about:

-template <> struct format_specifier<half    > { static constexpr const CUarray_format value = CU_AD_FORMAT_HALF;           };
+template <> struct format_specifier<__nv_half> { static constexpr const CUarray_format value = CU_AD_FORMAT_HALF;           };

?

[georgelyu]

The core problem is that half type is defined in cuda_fp16.hpp with a typedef __half half. Then when there is another math library that wants to typedef a half type, then there is conflict due to redefinition.

So, answering to your earlier questions, I think the only way to turn off this is to define CUDA_NO_HALF, because using __nv_half and other macros does not disable typedef __half half itself.

But I think another solution could be that we do not employ CUDA_NO_HALF, instead we define a new macro to control whether to include cuda_fp16.h in array.hpp. Because as you said, I'm not writing codes about half-precision either, but when I write #include <cuda/api.hpp>, then cuda_fp16.h is implicitly included, which could lead to potential conflicts. I think there should be a switch for people that don't want this part of features.

[eyalroz]

Well, you could still use CUDA_NO_HALF to control whether the aliasing of half occurs.

... on the other hand, I don't use cuda_fp16 for anything else, so... maybe let's go with your initial suggestion, but also, but CUDA_NO_HALF is specified, we don't including the cuda_fp16 header either. How about that?

[georgelyu]

I think that's good solution and I amended the commit. Modify anything if you'd like :)

[eyalroz]

Done, unless someone else complains and wants a more elaborate solution.

By the way - if you could take the time to give feedback on the changes since 0.5.6 on the development branch, that would help, since all I need to release 0.6 is either a bit longer to wait or more feedback about it from users, to give me enough confidence.

[georgelyu]

Cool. I'll try that and if I find any bug I'll submit an issue.

[eyalroz]

@georgelyu ... and if you don't - please email me