[CUDA] Fix order of vectorized ldg intrinsics' elements.

[jlebar]

Summary

The order is [x, y, z, w], not [w, x, y, z].

cfe/trunk/lib/Headers/__clang_cuda_intrinsics.h

inline __device__ char2 __ldg(const char2 *ptr) {
   ret.x = rv[0];
   ret.y = rv[1];
   return ret;
 }
 inline __device__ char4 __ldg(const char4 *ptr) {
   typedef char c4 __attribute__((ext_vector_type(4)));
   c4 rv = __nvvm_ldg_c4(reinterpret_cast<const c4 *>(ptr));
   char4 ret;
-  ret.w = rv[0];
-  ret.x = rv[1];
-  ret.y = rv[2];
-  ret.z = rv[3];
+  ret.x = rv[0];
+  ret.y = rv[1];
+  ret.z = rv[2];
+  ret.w = rv[3];
   return ret;
 }
 inline __device__ short2 __ldg(const short2 *ptr) {
   typedef short s2 __attribute__((ext_vector_type(2)));
   s2 rv = __nvvm_ldg_s2(reinterpret_cast<const s2 *>(ptr));
   short2 ret;
   ret.x = rv[0];
   ret.y = rv[1];
   return ret;
 }
 inline __device__ short4 __ldg(const short4 *ptr) {
   typedef short s4 __attribute__((ext_vector_type(4)));
   s4 rv = __nvvm_ldg_s4(reinterpret_cast<const s4 *>(ptr));
   short4 ret;
-  ret.w = rv[0];
-  ret.x = rv[1];
-  ret.y = rv[2];
-  ret.z = rv[3];
+  ret.x = rv[0];
+  ret.y = rv[1];
+  ret.z = rv[2];
+  ret.w = rv[3];
   return ret;
 }
 inline __device__ int2 __ldg(const int2 *ptr) {
   typedef int i2 __attribute__((ext_vector_type(2)));
   i2 rv = __nvvm_ldg_i2(reinterpret_cast<const i2 *>(ptr));
   int2 ret;
   ret.x = rv[0];
   ret.y = rv[1];
   return ret;
 }
 inline __device__ int4 __ldg(const int4 *ptr) {
   typedef int i4 __attribute__((ext_vector_type(4)));
   i4 rv = __nvvm_ldg_i4(reinterpret_cast<const i4 *>(ptr));
   int4 ret;
-  ret.w = rv[0];
-  ret.x = rv[1];
-  ret.y = rv[2];
-  ret.z = rv[3];
+  ret.x = rv[0];
+  ret.y = rv[1];
+  ret.z = rv[2];
+  ret.w = rv[3];
   return ret;
 }
 inline __device__ longlong2 __ldg(const longlong2 *ptr) {
   typedef long long ll2 __attribute__((ext_vector_type(2)));
   ll2 rv = __nvvm_ldg_ll2(reinterpret_cast<const ll2 *>(ptr));
   longlong2 ret;
   ret.x = rv[0];
   ret.y = rv[1];
   return ret;
 }
 
 inline __device__ uchar2 __ldg(const uchar2 *ptr) {
   typedef unsigned char uc2 __attribute__((ext_vector_type(2)));
   uc2 rv = __nvvm_ldg_uc2(reinterpret_cast<const uc2 *>(ptr));
   uchar2 ret;
   ret.x = rv[0];
   ret.y = rv[1];
   return ret;
 }
 inline __device__ uchar4 __ldg(const uchar4 *ptr) {
   typedef unsigned char uc4 __attribute__((ext_vector_type(4)));
   uc4 rv = __nvvm_ldg_uc4(reinterpret_cast<const uc4 *>(ptr));
   uchar4 ret;
-  ret.w = rv[0];
-  ret.x = rv[1];
-  ret.y = rv[2];
-  ret.z = rv[3];
+  ret.x = rv[0];
+  ret.y = rv[1];
+  ret.z = rv[2];
+  ret.w = rv[3];
   return ret;
 }
 inline __device__ ushort2 __ldg(const ushort2 *ptr) {
   typedef unsigned short us2 __attribute__((ext_vector_type(2)));
   us2 rv = __nvvm_ldg_us2(reinterpret_cast<const us2 *>(ptr));
   ushort2 ret;
   ret.x = rv[0];
   ret.y = rv[1];
   return ret;
 }
 inline __device__ ushort4 __ldg(const ushort4 *ptr) {
   typedef unsigned short us4 __attribute__((ext_vector_type(4)));
   us4 rv = __nvvm_ldg_us4(reinterpret_cast<const us4 *>(ptr));
   ushort4 ret;
-  ret.w = rv[0];
-  ret.x = rv[1];
-  ret.y = rv[2];
-  ret.z = rv[3];
+  ret.x = rv[0];
+  ret.y = rv[1];
+  ret.z = rv[2];
+  ret.w = rv[3];
   return ret;
 }
 inline __device__ uint2 __ldg(const uint2 *ptr) {
   typedef unsigned int ui2 __attribute__((ext_vector_type(2)));
   ui2 rv = __nvvm_ldg_ui2(reinterpret_cast<const ui2 *>(ptr));
   uint2 ret;
   ret.x = rv[0];
   ret.y = rv[1];
   return ret;
 }
 inline __device__ uint4 __ldg(const uint4 *ptr) {
   typedef unsigned int ui4 __attribute__((ext_vector_type(4)));
   ui4 rv = __nvvm_ldg_ui4(reinterpret_cast<const ui4 *>(ptr));
   uint4 ret;
-  ret.w = rv[0];
-  ret.x = rv[1];
-  ret.y = rv[2];
-  ret.z = rv[3];
+  ret.x = rv[0];
+  ret.y = rv[1];
+  ret.z = rv[2];
+  ret.w = rv[3];
   return ret;
 }
 inline __device__ ulonglong2 __ldg(const ulonglong2 *ptr) {
   typedef unsigned long long ull2 __attribute__((ext_vector_type(2)));
   ull2 rv = __nvvm_ldg_ull2(reinterpret_cast<const ull2 *>(ptr));
   ulonglong2 ret;
   ret.x = rv[0];
   ret.y = rv[1];
   return ret;
 }
 
 inline __device__ float2 __ldg(const float2 *ptr) {
   typedef float f2 __attribute__((ext_vector_type(2)));
   f2 rv = __nvvm_ldg_f2(reinterpret_cast<const f2 *>(ptr));
   float2 ret;
   ret.x = rv[0];
   ret.y = rv[1];
   return ret;
 }
 inline __device__ float4 __ldg(const float4 *ptr) {
   typedef float f4 __attribute__((ext_vector_type(4)));
   f4 rv = __nvvm_ldg_f4(reinterpret_cast<const f4 *>(ptr));
   float4 ret;
-  ret.w = rv[0];
-  ret.x = rv[1];
-  ret.y = rv[2];
-  ret.z = rv[3];
+  ret.x = rv[0];
+  ret.y = rv[1];
+  ret.z = rv[2];
+  ret.w = rv[3];
   return ret;
 }
 inline __device__ double2 __ldg(const double2 *ptr) {
   typedef double d2 __attribute__((ext_vector_type(2)));
   d2 rv = __nvvm_ldg_d2(reinterpret_cast<const d2 *>(ptr));
   double2 ret;
   ret.x = rv[0];
   ret.y = rv[1];