From 1fcde52afdf9a54f53265effb9beb2835e25879a Mon Sep 17 00:00:00 2001 From: KlausT Date: Sat, 22 Nov 2014 21:12:19 +0100 Subject: [PATCH] small groestl optimizations --- bitslice_transformations_quad.cu | 47 ++++++++++++++++---------------- 1 file changed, 24 insertions(+), 23 deletions(-) diff --git a/bitslice_transformations_quad.cu b/bitslice_transformations_quad.cu index 48c7f04..b5bfb9c 100644 --- a/bitslice_transformations_quad.cu +++ b/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; } }