X-Git-Url: https://www.flypig.org.uk/git/?a=blobdiff_plain;f=OpenCL%2Famp_a1.cl;h=d18e37c971ae2163a7856356d84bda554e1d9ad5;hb=161a6eb4bc643d8e636e96eda613f5137d30da59;hp=69b9898c47958a83d131e8a1676ae647b27853b3;hpb=bacc1049e3152771d306d97ce858f05d4024d331;p=hashcat.git

diff --git a/OpenCL/amp_a1.cl b/OpenCL/amp_a1.cl
index 69b9898..d18e37c 100644
--- a/OpenCL/amp_a1.cl
+++ b/OpenCL/amp_a1.cl
@@ -3,13 +3,13 @@
  * License.....: MIT
  */
 
-#include "include/constants.h"
-#include "include/kernel_vendor.h"
-#include "OpenCL/types_ocl.c"
+#include "inc_hash_constants.h"
+#include "inc_vendor.cl"
+#include "inc_types.cl"
 
-static void switch_buffer_by_offset (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 offset)
+inline void switch_buffer_by_offset_le (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 offset)
 {
-  #if defined IS_AMD || defined IS_UNKNOWN
+  #if defined IS_AMD || defined IS_GENERIC
   const int offset_mod_4 = offset & 3;
 
   const int offset_minus_4 = 4 - offset;
@@ -721,7 +721,7 @@ static void switch_buffer_by_offset (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4],
   #endif
 }
 
-__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) amp (__global pw_t *pws, __global pw_t *pws_amp, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, const u32 combs_mode, const u32 gid_max)
+__kernel void amp (__global pw_t *pws, __global pw_t *pws_amp, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, const u32 combs_mode, const u32 gid_max)
 {
   const u32 gid = get_global_id (0);
 
@@ -789,12 +789,12 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) amp (__global pw_
 
   if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
   {
-    switch_buffer_by_offset (wordr0, wordr1, wordr2, wordr3, pw_l_len);
+    switch_buffer_by_offset_le (wordr0, wordr1, wordr2, wordr3, pw_l_len);
   }
 
   if (combs_mode == COMBINATOR_MODE_BASE_RIGHT)
   {
-    switch_buffer_by_offset (wordl0, wordl1, wordl2, wordl3, pw_r_len);
+    switch_buffer_by_offset_le (wordl0, wordl1, wordl2, wordl3, pw_r_len);
   }
 
   u32 w0[4];