From 55f653f37464dfbe1c8ed2f1e57da96641d96b46 Mon Sep 17 00:00:00 2001 From: jsteube Date: Fri, 8 Sep 2017 19:47:56 +0200 Subject: [PATCH] Get rid of volatile in TrueCrypt kernels --- OpenCL/m06212.cl | 16 ---------------- OpenCL/m06213.cl | 24 ------------------------ OpenCL/m06222.cl | 16 ---------------- OpenCL/m06223.cl | 28 ++-------------------------- OpenCL/m06232.cl | 16 ---------------- OpenCL/m06233.cl | 24 ------------------------ 6 files changed, 2 insertions(+), 122 deletions(-) diff --git a/OpenCL/m06212.cl b/OpenCL/m06212.cl index 0762699e8..61351ef4f 100644 --- a/OpenCL/m06212.cl +++ b/OpenCL/m06212.cl @@ -331,11 +331,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (gid >= gid_max) return; - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey1[8]; - #else u32 ukey1[8]; - #endif ukey1[0] = tmps[gid].out[ 0]; ukey1[1] = tmps[gid].out[ 1]; @@ -346,11 +342,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global const kernel_rule_t *rul ukey1[6] = tmps[gid].out[ 6]; ukey1[7] = tmps[gid].out[ 7]; - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey2[8]; - #else u32 ukey2[8]; - #endif ukey2[0] = tmps[gid].out[ 8]; ukey2[1] = tmps[gid].out[ 9]; @@ -385,11 +377,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global const kernel_rule_t *rul } } - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey3[8]; - #else u32 ukey3[8]; - #endif ukey3[0] = tmps[gid].out[16]; ukey3[1] = tmps[gid].out[17]; @@ -400,11 +388,7 @@ __kernel void m06212_comp (__global pw_t *pws, __global const kernel_rule_t *rul ukey3[6] = tmps[gid].out[22]; ukey3[7] = tmps[gid].out[23]; - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey4[8]; - #else u32 ukey4[8]; - #endif ukey4[0] = tmps[gid].out[24]; ukey4[1] = tmps[gid].out[25]; diff --git a/OpenCL/m06213.cl b/OpenCL/m06213.cl index 3aa2116dc..0ee3441b8 100644 --- a/OpenCL/m06213.cl +++ b/OpenCL/m06213.cl @@ -331,11 +331,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (gid >= gid_max) return; - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey1[8]; - #else u32 ukey1[8]; - #endif ukey1[0] = tmps[gid].out[ 0]; ukey1[1] = tmps[gid].out[ 1]; @@ -346,11 +342,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul ukey1[6] = tmps[gid].out[ 6]; ukey1[7] = tmps[gid].out[ 7]; - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey2[8]; - #else u32 ukey2[8]; - #endif ukey2[0] = tmps[gid].out[ 8]; ukey2[1] = tmps[gid].out[ 9]; @@ -385,11 +377,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul } } - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey3[8]; - #else u32 ukey3[8]; - #endif ukey3[0] = tmps[gid].out[16]; ukey3[1] = tmps[gid].out[17]; @@ -400,11 +388,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul ukey3[6] = tmps[gid].out[22]; ukey3[7] = tmps[gid].out[23]; - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey4[8]; - #else u32 ukey4[8]; - #endif ukey4[0] = tmps[gid].out[24]; ukey4[1] = tmps[gid].out[25]; @@ -439,11 +423,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul } } - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey5[8]; - #else u32 ukey5[8]; - #endif ukey5[0] = tmps[gid].out[32]; ukey5[1] = tmps[gid].out[33]; @@ -454,11 +434,7 @@ __kernel void m06213_comp (__global pw_t *pws, __global const kernel_rule_t *rul ukey5[6] = tmps[gid].out[38]; ukey5[7] = tmps[gid].out[39]; - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey6[8]; - #else u32 ukey6[8]; - #endif ukey6[0] = tmps[gid].out[40]; ukey6[1] = tmps[gid].out[41]; diff --git a/OpenCL/m06222.cl b/OpenCL/m06222.cl index cb533cb16..1f2f40030 100644 --- a/OpenCL/m06222.cl +++ b/OpenCL/m06222.cl @@ -475,11 +475,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (gid >= gid_max) return; - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey1[8]; - #else u32 ukey1[8]; - #endif ukey1[0] = swap32_S (h32_from_64_S (tmps[gid].out[0])); ukey1[1] = swap32_S (l32_from_64_S (tmps[gid].out[0])); @@ -490,11 +486,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global const kernel_rule_t *rul ukey1[6] = swap32_S (h32_from_64_S (tmps[gid].out[3])); ukey1[7] = swap32_S (l32_from_64_S (tmps[gid].out[3])); - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey2[8]; - #else u32 ukey2[8]; - #endif ukey2[0] = swap32_S (h32_from_64_S (tmps[gid].out[4])); ukey2[1] = swap32_S (l32_from_64_S (tmps[gid].out[4])); @@ -529,11 +521,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global const kernel_rule_t *rul } } - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey3[8]; - #else u32 ukey3[8]; - #endif ukey3[0] = swap32_S (h32_from_64_S (tmps[gid].out[ 8])); ukey3[1] = swap32_S (l32_from_64_S (tmps[gid].out[ 8])); @@ -544,11 +532,7 @@ __kernel void m06222_comp (__global pw_t *pws, __global const kernel_rule_t *rul ukey3[6] = swap32_S (h32_from_64_S (tmps[gid].out[11])); ukey3[7] = swap32_S (l32_from_64_S (tmps[gid].out[11])); - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey4[8]; - #else u32 ukey4[8]; - #endif ukey4[0] = swap32_S (h32_from_64_S (tmps[gid].out[12])); ukey4[1] = swap32_S (l32_from_64_S (tmps[gid].out[12])); diff --git a/OpenCL/m06223.cl b/OpenCL/m06223.cl index 662f0a84e..e540d2ec2 100644 --- a/OpenCL/m06223.cl +++ b/OpenCL/m06223.cl @@ -475,11 +475,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (gid >= gid_max) return; - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey1[8]; - #else u32 ukey1[8]; - #endif ukey1[0] = swap32_S (h32_from_64_S (tmps[gid].out[0])); ukey1[1] = swap32_S (l32_from_64_S (tmps[gid].out[0])); @@ -490,11 +486,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul ukey1[6] = swap32_S (h32_from_64_S (tmps[gid].out[3])); ukey1[7] = swap32_S (l32_from_64_S (tmps[gid].out[3])); - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey2[8]; - #else u32 ukey2[8]; - #endif ukey2[0] = swap32_S (h32_from_64_S (tmps[gid].out[4])); ukey2[1] = swap32_S (l32_from_64_S (tmps[gid].out[4])); @@ -529,11 +521,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul } } - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey3[8]; - #else u32 ukey3[8]; - #endif ukey3[0] = swap32_S (h32_from_64_S (tmps[gid].out[ 8])); ukey3[1] = swap32_S (l32_from_64_S (tmps[gid].out[ 8])); @@ -544,11 +532,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul ukey3[6] = swap32_S (h32_from_64_S (tmps[gid].out[11])); ukey3[7] = swap32_S (l32_from_64_S (tmps[gid].out[11])); - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey4[8]; - #else u32 ukey4[8]; - #endif ukey4[0] = swap32_S (h32_from_64_S (tmps[gid].out[12])); ukey4[1] = swap32_S (l32_from_64_S (tmps[gid].out[12])); @@ -583,11 +567,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul } } - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey5[8]; - #else - volatile u32 ukey5[8]; - #endif + u32 ukey5[8]; ukey5[0] = swap32_S (h32_from_64_S (tmps[gid].out[16])); ukey5[1] = swap32_S (l32_from_64_S (tmps[gid].out[16])); @@ -598,11 +578,7 @@ __kernel void m06223_comp (__global pw_t *pws, __global const kernel_rule_t *rul ukey5[6] = swap32_S (h32_from_64_S (tmps[gid].out[19])); ukey5[7] = swap32_S (l32_from_64_S (tmps[gid].out[19])); - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey6[8]; - #else - volatile u32 ukey6[8]; - #endif + u32 ukey6[8]; ukey6[0] = swap32_S (h32_from_64_S (tmps[gid].out[20])); ukey6[1] = swap32_S (l32_from_64_S (tmps[gid].out[20])); diff --git a/OpenCL/m06232.cl b/OpenCL/m06232.cl index eaac27c99..e2e2a1a3e 100644 --- a/OpenCL/m06232.cl +++ b/OpenCL/m06232.cl @@ -615,11 +615,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (gid >= gid_max) return; - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey1[8]; - #else u32 ukey1[8]; - #endif ukey1[0] = swap32_S (tmps[gid].out[ 0]); ukey1[1] = swap32_S (tmps[gid].out[ 1]); @@ -630,11 +626,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global const kernel_rule_t *rul ukey1[6] = swap32_S (tmps[gid].out[ 6]); ukey1[7] = swap32_S (tmps[gid].out[ 7]); - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey2[8]; - #else u32 ukey2[8]; - #endif ukey2[0] = swap32_S (tmps[gid].out[ 8]); ukey2[1] = swap32_S (tmps[gid].out[ 9]); @@ -669,11 +661,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global const kernel_rule_t *rul } } - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey3[8]; - #else u32 ukey3[8]; - #endif ukey3[0] = swap32_S (tmps[gid].out[16]); ukey3[1] = swap32_S (tmps[gid].out[17]); @@ -684,11 +672,7 @@ __kernel void m06232_comp (__global pw_t *pws, __global const kernel_rule_t *rul ukey3[6] = swap32_S (tmps[gid].out[22]); ukey3[7] = swap32_S (tmps[gid].out[23]); - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey4[8]; - #else u32 ukey4[8]; - #endif ukey4[0] = swap32_S (tmps[gid].out[24]); ukey4[1] = swap32_S (tmps[gid].out[25]); diff --git a/OpenCL/m06233.cl b/OpenCL/m06233.cl index bcf479528..601b0fb45 100644 --- a/OpenCL/m06233.cl +++ b/OpenCL/m06233.cl @@ -615,11 +615,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul if (gid >= gid_max) return; - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey1[8]; - #else u32 ukey1[8]; - #endif ukey1[0] = swap32_S (tmps[gid].out[ 0]); ukey1[1] = swap32_S (tmps[gid].out[ 1]); @@ -630,11 +626,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul ukey1[6] = swap32_S (tmps[gid].out[ 6]); ukey1[7] = swap32_S (tmps[gid].out[ 7]); - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey2[8]; - #else u32 ukey2[8]; - #endif ukey2[0] = swap32_S (tmps[gid].out[ 8]); ukey2[1] = swap32_S (tmps[gid].out[ 9]); @@ -669,11 +661,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul } } - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey3[8]; - #else u32 ukey3[8]; - #endif ukey3[0] = swap32_S (tmps[gid].out[16]); ukey3[1] = swap32_S (tmps[gid].out[17]); @@ -684,11 +672,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul ukey3[6] = swap32_S (tmps[gid].out[22]); ukey3[7] = swap32_S (tmps[gid].out[23]); - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey4[8]; - #else u32 ukey4[8]; - #endif ukey4[0] = swap32_S (tmps[gid].out[24]); ukey4[1] = swap32_S (tmps[gid].out[25]); @@ -723,11 +707,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul } } - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey5[8]; - #else u32 ukey5[8]; - #endif ukey5[0] = swap32_S (tmps[gid].out[32]); ukey5[1] = swap32_S (tmps[gid].out[33]); @@ -738,11 +718,7 @@ __kernel void m06233_comp (__global pw_t *pws, __global const kernel_rule_t *rul ukey5[6] = swap32_S (tmps[gid].out[38]); ukey5[7] = swap32_S (tmps[gid].out[39]); - #if defined (IS_APPLE) && defined (IS_GPU) - volatile u32 ukey6[8]; - #else u32 ukey6[8]; - #endif ukey6[0] = swap32_S (tmps[gid].out[40]); ukey6[1] = swap32_S (tmps[gid].out[41]);