From 69ed4008b73b98a62067fa5cb7a371e1fd1fc22a Mon Sep 17 00:00:00 2001
From: Jared Hoberock <jaredhoberock@gmail.com>
Date: Fri, 26 Apr 2013 10:54:06 -0700
Subject: [PATCH] Check for signedness of char when performing key conversion
 in CUDA's radix sort.

---
 .../detail/b40c/radixsort_key_conversion.h    | 28 ++++++++++++-------
 1 file changed, 18 insertions(+), 10 deletions(-)

diff --git a/thrust/system/cuda/detail/detail/b40c/radixsort_key_conversion.h b/thrust/system/cuda/detail/detail/b40c/radixsort_key_conversion.h
index c1c5ad4db..a170f95e6 100644
--- a/thrust/system/cuda/detail/detail/b40c/radixsort_key_conversion.h
+++ b/thrust/system/cuda/detail/detail/b40c/radixsort_key_conversion.h
@@ -164,25 +164,33 @@ struct PostprocessKeyFunctor<double> {
 //
 
 template <> struct KeyConversion<char> {
-	typedef unsigned char UnsignedBits;
+  typedef unsigned char UnsignedBits;
 };
 
 template <>
 struct PreprocessKeyFunctor<char> {
-	__device__ __host__ __forceinline__ void operator()(unsigned char &converted_key) {
-		const unsigned int SIGN_MASK = 1u << ((sizeof(char) * 8) - 1);
-		converted_key ^= SIGN_MASK;	
-	}
-	__device__ __host__ __forceinline__ static bool MustApply(){ return true;}
+  __device__ __host__ __forceinline__ void operator()(unsigned char &converted_key) {
+    // char is unsigned on some platforms, so we have to check
+    if(std::numeric_limits<char>::is_signed)
+    {
+      const unsigned int SIGN_MASK = 1u << ((sizeof(char) * 8) - 1);
+      converted_key ^= SIGN_MASK;	
+    }
+  }
+  __device__ __host__ __forceinline__ static bool MustApply(){ return std::numeric_limits<char>::is_signed;}
 };
 
 template <>
 struct PostprocessKeyFunctor<char> {
-	__device__ __host__ __forceinline__ void operator()(unsigned char &converted_key)  {
-		const unsigned int SIGN_MASK = 1u << ((sizeof(char) * 8) - 1);
-		converted_key ^= SIGN_MASK;	
+  __device__ __host__ __forceinline__ void operator()(unsigned char &converted_key)  {
+    // char is unsigned on some platforms, so we have to check
+    if(std::numeric_limits<char>::is_signed)
+    {
+      const unsigned int SIGN_MASK = 1u << ((sizeof(char) * 8) - 1);
+      converted_key ^= SIGN_MASK;	
     }
-	__device__ __host__ __forceinline__ static bool MustApply(){ return true;}
+  }
+  __device__ __host__ __forceinline__ static bool MustApply(){ return std::numeric_limits<char>::is_signed;}
 };