More CUDA special backports to HIP

OpenCL/inc_platform.cl

@@ -309,8 +309,9 @@ DECLSPEC u64 rotr64_S (const u64 a, const int n)
   return out.v64;
 }
 
-#define FIXED_THREAD_COUNT(n) __launch_bounds__((n), 0)
-#define SYNC_THREADS() __syncthreads ()
+#define FIXED_THREAD_COUNT(n) __attribute__((amdgpu_flat_work_group_size (1, (n))))
+
+#define SYNC_THREADS() __builtin_amdgcn_s_barrier ()
 #endif
 
 #ifdef IS_OPENCL

OpenCL/m02500-pure.cl

@@ -775,11 +775,7 @@ KERNEL_FQ void m02500_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_eapol_t)
     s_te4[i] = te4[i];
   }
 
-  #if defined IS_CUDA || defined IS_HIP
-  __syncthreads();
-  #else
   SYNC_THREADS ();
-  #endif
 
   #else
 

OpenCL/m22000-pure.cl

@@ -797,11 +797,7 @@ KERNEL_FQ void m22000_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_t))
     s_te4[i] = te4[i];
   }
 
-  #if defined IS_CUDA || defined IS_HIP
-  __syncthreads();
-  #else
   SYNC_THREADS ();
-  #endif
 
   #else
 

OpenCL/m22001-pure.cl

@@ -610,11 +610,7 @@ KERNEL_FQ void m22001_aux3 (KERN_ATTR_TMPS_ESALT (wpa_pmk_tmp_t, wpa_t))
     s_te4[i] = te4[i];
   }
 
-  #if defined IS_CUDA || defined IS_HIP
-  __syncthreads();
-  #else
   SYNC_THREADS ();
-  #endif
 
   #else