Fix integer overflow in CUDA decoder

src/cuda_zfp/shared.h

@@ -223,29 +223,40 @@ __device__
 static void
 inv_lift(Int* p)
 {
-	Int x, y, z, w;
-	x = *p; p += s;
-	y = *p; p += s;
-	z = *p; p += s;
-	w = *p; p += s;
-
-	/*
-	** non-orthogonal transform
-	**       ( 4  6 -4 -1) (x)
-	** 1/4 * ( 4  2  4  5) (y)
-	**       ( 4 -2  4 -5) (z)
-	**       ( 4 -6 -4  1) (w)
-	*/
-	y += w >> 1; w -= y >> 1;
-	y += w; w <<= 1; w -= y;
-	z += x; x <<= 1; x -= z;
-	y += z; z <<= 1; z -= y;
-	w += x; x <<= 1; x -= w;
-
-	p -= s; *p = w;
-	p -= s; *p = z;
-	p -= s; *p = y;
-	p -= s; *p = x;
+  Int x, y, z, w;
+  x = *p; p += s;
+  y = *p; p += s;
+  z = *p; p += s;
+  w = *p; p += s;
+
+  /*
+  ** non-orthogonal transform
+  **
+  **       ( 4  6 -4 -1) (x)
+  ** 1/4 * ( 4  2  4  5) (y)
+  **       ( 4 -2  4 -5) (z)
+  **       ( 4 -6 -4  1) (w)
+  **
+  ** original lifted version, which invokes UB due to signed left shift and
+  ** integer overflow:
+  **
+  ** y += w >> 1; w -= y >> 1;
+  ** y += w; w <<= 1; w -= y;
+  ** z += x; x <<= 1; x -= z;
+  ** y += z; z <<= 1; z -= y;
+  ** w += x; x <<= 1; x -= w;
+  */
+
+  y += w >> 1; w -= y >> 1;
+  y += w; w -= y - w;
+  z += x; x -= z - x;
+  y += z; z -= y - z;
+  w += x; x -= w - x;
+
+  p -= s; *p = w;
+  p -= s; *p = z;
+  p -= s; *p = y;
+  p -= s; *p = x;
 }