small groestl optimizations

bitslice_transformations_quad.cu

@@ -25,10 +25,9 @@ void to_bitslice_quad(uint32_t *input, uint32_t *output)
             input[i] = __byte_perm(input[i], 0, 0x1032);
             other[i] = __byte_perm(other[i], 0, 0x1032);
         }
-        output[i] = 0;
     }
 
-    output[ 0] |=  (input[ 0] & 0x00000001);
+    output[ 0] =  (input[ 0] & 0x00000001);
     output[ 0] |= ((other[ 0] & 0x00000001) << 1);
     output[ 0] |= ((input[ 1] & 0x00000001) << 2);
     output[ 0] |= ((other[ 1] & 0x00000001) << 3);
@@ -61,7 +60,7 @@ void to_bitslice_quad(uint32_t *input, uint32_t *output)
     output[ 0] |= ((input[ 7] & 0x00000100) <<22);
     output[ 0] |= ((other[ 7] & 0x00000100) <<23);
 
-    output[ 1] |= ((input[ 0] & 0x00000002) >> 1);
+    output[ 1] = ((input[ 0] & 0x00000002) >> 1);
     output[ 1] |=  (other[ 0] & 0x00000002);
     output[ 1] |= ((input[ 1] & 0x00000002) << 1);
     output[ 1] |= ((other[ 1] & 0x00000002) << 2);
@@ -94,7 +93,7 @@ void to_bitslice_quad(uint32_t *input, uint32_t *output)
     output[ 1] |= ((input[ 7] & 0x00000200) <<21);
     output[ 1] |= ((other[ 7] & 0x00000200) <<22);
 
-    output[ 2] |= ((input[ 0] & 0x00000004) >> 2);
+    output[ 2] = ((input[ 0] & 0x00000004) >> 2);
     output[ 2] |= ((other[ 0] & 0x00000004) >> 1);
     output[ 2] |=  (input[ 1] & 0x00000004);
     output[ 2] |= ((other[ 1] & 0x00000004) << 1);
@@ -126,7 +125,8 @@ void to_bitslice_quad(uint32_t *input, uint32_t *output)
     output[ 2] |= ((other[ 6] & 0x00000400) <<19);
     output[ 2] |= ((input[ 7] & 0x00000400) <<20);
     output[ 2] |= ((other[ 7] & 0x00000400) <<21);
-    output[ 3] |= ((input[ 0] & 0x00000008) >> 3);
+
+    output[ 3] = ((input[ 0] & 0x00000008) >> 3);
     output[ 3] |= ((other[ 0] & 0x00000008) >> 2);
     output[ 3] |= ((input[ 1] & 0x00000008) >> 1);
     output[ 3] |=  (other[ 1] & 0x00000008);
@@ -158,7 +158,8 @@ void to_bitslice_quad(uint32_t *input, uint32_t *output)
     output[ 3] |= ((other[ 6] & 0x00000800) <<18);
     output[ 3] |= ((input[ 7] & 0x00000800) <<19);
     output[ 3] |= ((other[ 7] & 0x00000800) <<20);
-    output[ 4] |= ((input[ 0] & 0x00000010) >> 4);
+
+    output[ 4] = ((input[ 0] & 0x00000010) >> 4);
     output[ 4] |= ((other[ 0] & 0x00000010) >> 3);
     output[ 4] |= ((input[ 1] & 0x00000010) >> 2);
     output[ 4] |= ((other[ 1] & 0x00000010) >> 1);
@@ -190,7 +191,8 @@ void to_bitslice_quad(uint32_t *input, uint32_t *output)
     output[ 4] |= ((other[ 6] & 0x00001000) <<17);
     output[ 4] |= ((input[ 7] & 0x00001000) <<18);
     output[ 4] |= ((other[ 7] & 0x00001000) <<19);
-    output[ 5] |= ((input[ 0] & 0x00000020) >> 5);
+
+    output[ 5] = ((input[ 0] & 0x00000020) >> 5);
     output[ 5] |= ((other[ 0] & 0x00000020) >> 4);
     output[ 5] |= ((input[ 1] & 0x00000020) >> 3);
     output[ 5] |= ((other[ 1] & 0x00000020) >> 2);
@@ -222,7 +224,8 @@ void to_bitslice_quad(uint32_t *input, uint32_t *output)
     output[ 5] |= ((other[ 6] & 0x00002000) <<16);
     output[ 5] |= ((input[ 7] & 0x00002000) <<17);
     output[ 5] |= ((other[ 7] & 0x00002000) <<18);
-    output[ 6] |= ((input[ 0] & 0x00000040) >> 6);
+
+    output[ 6] = ((input[ 0] & 0x00000040) >> 6);
     output[ 6] |= ((other[ 0] & 0x00000040) >> 5);
     output[ 6] |= ((input[ 1] & 0x00000040) >> 4);
     output[ 6] |= ((other[ 1] & 0x00000040) >> 3);
@@ -254,7 +257,8 @@ void to_bitslice_quad(uint32_t *input, uint32_t *output)
     output[ 6] |= ((other[ 6] & 0x00004000) <<15);
     output[ 6] |= ((input[ 7] & 0x00004000) <<16);
     output[ 6] |= ((other[ 7] & 0x00004000) <<17);
-    output[ 7] |= ((input[ 0] & 0x00000080) >> 7);
+
+    output[ 7] = ((input[ 0] & 0x00000080) >> 7);
     output[ 7] |= ((other[ 0] & 0x00000080) >> 6);
     output[ 7] |= ((input[ 1] & 0x00000080) >> 5);
     output[ 7] |= ((other[ 1] & 0x00000080) >> 4);
@@ -291,10 +295,7 @@ void to_bitslice_quad(uint32_t *input, uint32_t *output)
 __device__ __forceinline__
 void from_bitslice_quad(uint32_t *input, uint32_t *output)
 {
-#pragma unroll 8
-    for (int i=0; i < 16; i+=2) output[i] = 0;
-
-    output[ 0] |= ((input[ 0] & 0x00000100) >> 8);
+    output[ 0] = ((input[ 0] & 0x00000100) >> 8);
     output[ 0] |= ((input[ 1] & 0x00000100) >> 7);
     output[ 0] |= ((input[ 2] & 0x00000100) >> 6);
     output[ 0] |= ((input[ 3] & 0x00000100) >> 5);
@@ -310,7 +311,7 @@ void from_bitslice_quad(uint32_t *input, uint32_t *output)
     output[ 0] |= ((input[ 5] & 0x01000000) >>11);
     output[ 0] |= ((input[ 6] & 0x01000000) >>10);
     output[ 0] |= ((input[ 7] & 0x01000000) >> 9);
-    output[ 2] |= ((input[ 0] & 0x00000200) >> 9);
+    output[ 2] = ((input[ 0] & 0x00000200) >> 9);
     output[ 2] |= ((input[ 1] & 0x00000200) >> 8);
     output[ 2] |= ((input[ 2] & 0x00000200) >> 7);
     output[ 2] |= ((input[ 3] & 0x00000200) >> 6);
@@ -326,7 +327,7 @@ void from_bitslice_quad(uint32_t *input, uint32_t *output)
     output[ 2] |= ((input[ 5] & 0x02000000) >>12);
     output[ 2] |= ((input[ 6] & 0x02000000) >>11);
     output[ 2] |= ((input[ 7] & 0x02000000) >>10);
-    output[ 4] |= ((input[ 0] & 0x00000400) >>10);
+    output[ 4] = ((input[ 0] & 0x00000400) >>10);
     output[ 4] |= ((input[ 1] & 0x00000400) >> 9);
     output[ 4] |= ((input[ 2] & 0x00000400) >> 8);
     output[ 4] |= ((input[ 3] & 0x00000400) >> 7);
@@ -342,7 +343,7 @@ void from_bitslice_quad(uint32_t *input, uint32_t *output)
     output[ 4] |= ((input[ 5] & 0x04000000) >>13);
     output[ 4] |= ((input[ 6] & 0x04000000) >>12);
     output[ 4] |= ((input[ 7] & 0x04000000) >>11);
-    output[ 6] |= ((input[ 0] & 0x00000800) >>11);
+    output[ 6] = ((input[ 0] & 0x00000800) >>11);
     output[ 6] |= ((input[ 1] & 0x00000800) >>10);
     output[ 6] |= ((input[ 2] & 0x00000800) >> 9);
     output[ 6] |= ((input[ 3] & 0x00000800) >> 8);
@@ -358,7 +359,7 @@ void from_bitslice_quad(uint32_t *input, uint32_t *output)
     output[ 6] |= ((input[ 5] & 0x08000000) >>14);
     output[ 6] |= ((input[ 6] & 0x08000000) >>13);
     output[ 6] |= ((input[ 7] & 0x08000000) >>12);
-    output[ 8] |= ((input[ 0] & 0x00001000) >>12);
+    output[ 8] = ((input[ 0] & 0x00001000) >>12);
     output[ 8] |= ((input[ 1] & 0x00001000) >>11);
     output[ 8] |= ((input[ 2] & 0x00001000) >>10);
     output[ 8] |= ((input[ 3] & 0x00001000) >> 9);
@@ -374,7 +375,7 @@ void from_bitslice_quad(uint32_t *input, uint32_t *output)
     output[ 8] |= ((input[ 5] & 0x10000000) >>15);
     output[ 8] |= ((input[ 6] & 0x10000000) >>14);
     output[ 8] |= ((input[ 7] & 0x10000000) >>13);
-    output[10] |= ((input[ 0] & 0x00002000) >>13);
+    output[10] = ((input[ 0] & 0x00002000) >>13);
     output[10] |= ((input[ 1] & 0x00002000) >>12);
     output[10] |= ((input[ 2] & 0x00002000) >>11);
     output[10] |= ((input[ 3] & 0x00002000) >>10);
@@ -390,7 +391,7 @@ void from_bitslice_quad(uint32_t *input, uint32_t *output)
     output[10] |= ((input[ 5] & 0x20000000) >>16);
     output[10] |= ((input[ 6] & 0x20000000) >>15);
     output[10] |= ((input[ 7] & 0x20000000) >>14);
-    output[12] |= ((input[ 0] & 0x00004000) >>14);
+    output[12] = ((input[ 0] & 0x00004000) >>14);
     output[12] |= ((input[ 1] & 0x00004000) >>13);
     output[12] |= ((input[ 2] & 0x00004000) >>12);
     output[12] |= ((input[ 3] & 0x00004000) >>11);
@@ -406,7 +407,7 @@ void from_bitslice_quad(uint32_t *input, uint32_t *output)
     output[12] |= ((input[ 5] & 0x40000000) >>17);
     output[12] |= ((input[ 6] & 0x40000000) >>16);
     output[12] |= ((input[ 7] & 0x40000000) >>15);
-    output[14] |= ((input[ 0] & 0x00008000) >>15);
+    output[14] = ((input[ 0] & 0x00008000) >>15);
     output[14] |= ((input[ 1] & 0x00008000) >>14);
     output[14] |= ((input[ 2] & 0x00008000) >>13);
     output[14] |= ((input[ 3] & 0x00008000) >>12);
@@ -426,9 +427,9 @@ void from_bitslice_quad(uint32_t *input, uint32_t *output)
 #pragma unroll 8
     for (int i = 0; i < 16; i+=2) {
         if (threadIdx.x & 1) output[i] = __byte_perm(output[i], 0, 0x1032);
-        output[i] = __byte_perm(output[i], __shfl((int)output[i], (threadIdx.x+1)%4, 4), 0x7610);
-        output[i+1] = __shfl((int)output[i], (threadIdx.x+2)%4, 4);
-        if (threadIdx.x % 4) output[i] = output[i+1] = 0;
+        output[i] = __byte_perm(output[i], __shfl((int)output[i], (threadIdx.x+1)&3, 4), 0x7610);
+        output[i+1] = __shfl((int)output[i], (threadIdx.x+2)&3, 4);
+        if (threadIdx.x & 3) output[i] = output[i+1] = 0;
     }
 }