From 7645a6f00fcffde9a3b479c2b1cbb50d18cdc9ca Mon Sep 17 00:00:00 2001 From: jsteube Date: Tue, 16 Feb 2016 13:24:36 +0100 Subject: [PATCH] Fix warp bug on -m 8500 --- OpenCL/m08500_a0.cl | 4 ++-- OpenCL/m08500_a1.cl | 4 ++-- OpenCL/m08500_a3.cl | 4 ++-- 3 files changed, 6 insertions(+), 6 deletions(-) diff --git a/OpenCL/m08500_a0.cl b/OpenCL/m08500_a0.cl index 15dc885dc..71bdc2805 100644 --- a/OpenCL/m08500_a0.cl +++ b/OpenCL/m08500_a0.cl @@ -626,7 +626,7 @@ __kernel void m08500_m04 (__global pw_t *pws, __global kernel_rule_t * rules_bu data[0] = salt_buf0[0]; data[1] = salt_buf0[1]; - u32x iv[2]; + volatile u32x iv[2]; _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); @@ -760,7 +760,7 @@ __kernel void m08500_s04 (__global pw_t *pws, __global kernel_rule_t * rules_bu data[0] = salt_buf0[0]; data[1] = salt_buf0[1]; - u32x iv[2]; + volatile u32x iv[2]; _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); diff --git a/OpenCL/m08500_a1.cl b/OpenCL/m08500_a1.cl index ccefd4790..f3f653d82 100644 --- a/OpenCL/m08500_a1.cl +++ b/OpenCL/m08500_a1.cl @@ -685,7 +685,7 @@ __kernel void m08500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, data[0] = salt_buf0[0]; data[1] = salt_buf0[1]; - u32 iv[2]; + volatile u32 iv[2]; _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); @@ -898,7 +898,7 @@ __kernel void m08500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, data[0] = salt_buf0[0]; data[1] = salt_buf0[1]; - u32 iv[2]; + volatile u32 iv[2]; _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); diff --git a/OpenCL/m08500_a3.cl b/OpenCL/m08500_a3.cl index ebc84702a..2e9b9c5be 100644 --- a/OpenCL/m08500_a3.cl +++ b/OpenCL/m08500_a3.cl @@ -570,7 +570,7 @@ static void m08500m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 data[0] = salt_buf0[0]; data[1] = salt_buf0[1]; - u32x iv[2]; + volatile u32x iv[2]; _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); @@ -642,7 +642,7 @@ static void m08500s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 data[0] = salt_buf0[0]; data[1] = salt_buf0[1]; - u32x iv[2]; + volatile u32x iv[2]; _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);