fixes to avoid crashing on gpu on multiple inner loops

OpenCL/inc_zip_inflate.cl

@@ -180,7 +180,7 @@ typedef int mz_bool;
 
 typedef mz_uint64 tinfl_bit_buf_t;
 
-void memcpy(void *dest, const void *src, size_t n){
+void memcpy(void *dest, const void *src, u32 n){
   char *csrc = (char *)src;
   char *cdest = (char *)dest;
   for (int i=0; i<n; i++){
@@ -189,7 +189,7 @@ void memcpy(void *dest, const void *src, size_t n){
 }
 
 
-void *memset(void *s, int c, u32 len){
+void *memset(u8 *s, int c, u32 len){
   u8 *dst = s;
   while (len > 0) {
       *dst = (u8) c;
@@ -206,7 +206,7 @@ void *memset(void *s, int c, u32 len){
 #define TINFL_LZ_DICT_SIZE 32768
 #define TINFL_MEMCPY(d, s, l) memcpy(d, s, l)
 #define TINFL_MEMCPY_G(d, s, l, p) memcpy_g(d, s, l, p)
-#define TINFL_MEMSET(p, c, l) memset(p, c, l)
+#define TINFL_MEMSET(p, c, l) memset(p, c, (u32)l)
 #define MZ_CLEAR_OBJ(obj) memset(&(obj), 0, sizeof(obj))
 
 #define TINFL_CR_FINISH }
@@ -418,7 +418,7 @@ typedef struct
 
 typedef struct mz_stream_s
 {
-    __global const unsigned char *next_in; /* pointer to next byte to read */
+    GLOBAL_AS const unsigned char *next_in; /* pointer to next byte to read */
     unsigned int avail_in;        /* number of bytes available at next_in */
     mz_ulong total_in;            /* total number of bytes consumed so far */
 
@@ -427,7 +427,7 @@ typedef struct mz_stream_s
     mz_ulong total_out;      /* total number of bytes produced so far */
 
     char *msg;                       /* error msg (unused) */
-    struct inflate_state *state; /* internal state, allocated by zalloc/zfree */
+    inflate_state *state; /* internal state, allocated by zalloc/zfree */
 
     void *opaque;         /* heap alloc function user pointer */
 
@@ -472,15 +472,15 @@ const mz_uint8 pIn_xor_byte (const mz_uint8 c, mz_streamp pStream)
 }
 
 
-void memcpy_g(void *dest, __global const void *src, size_t n, mz_streamp pStream){
+void memcpy_g(void *dest, GLOBAL_AS const void *src, size_t n, mz_streamp pStream){
-  __global char *csrc = (__global char *)src;
+  GLOBAL_AS char *csrc = (GLOBAL_AS char *)src;
   char *cdest = (char *)dest;
   for (int i=0; i<n; i++){
     cdest[i] = pIn_xor_byte (csrc[i], pStream);
   }
 }
 
-tinfl_status tinfl_decompress(tinfl_decompressor *r, __global const mz_uint8 *pIn_buf_next, size_t *pIn_buf_size, mz_uint8 *pOut_buf_start, mz_uint8 *pOut_buf_next, size_t *pOut_buf_size, const mz_uint32 decomp_flags, mz_streamp pStream)
+tinfl_status tinfl_decompress(tinfl_decompressor *r, GLOBAL_AS const mz_uint8 *pIn_buf_next, size_t *pIn_buf_size, mz_uint8 *pOut_buf_start, mz_uint8 *pOut_buf_next, size_t *pOut_buf_size, const mz_uint32 decomp_flags, mz_streamp pStream)
 {
 
     const int s_length_base[31] = { 3, 4, 5, 6, 7, 8, 9, 10, 11, 13, 15, 17, 19, 23, 27, 31, 35, 43, 51, 59, 67, 83, 99, 115, 131, 163, 195, 227, 258, 0, 0 };
@@ -493,8 +493,8 @@ tinfl_status tinfl_decompress(tinfl_decompressor *r, __global const mz_uint8 *pI
     tinfl_status status = TINFL_STATUS_FAILED;
     mz_uint32 num_bits, dist, counter, num_extra;
     tinfl_bit_buf_t bit_buf;
-    __global const mz_uint8 *pIn_buf_cur = pIn_buf_next;
+    GLOBAL_AS const mz_uint8 *pIn_buf_cur = pIn_buf_next;
-    __global const mz_uint8 *pIn_buf_end = pIn_buf_next + *pIn_buf_size;
+    GLOBAL_AS const mz_uint8 *pIn_buf_end = pIn_buf_next + *pIn_buf_size;
     mz_uint8 *pOut_buf_cur = pOut_buf_next, *const pOut_buf_end = pOut_buf_next + *pOut_buf_size;
     size_t out_buf_size_mask = (decomp_flags & TINFL_FLAG_USING_NON_WRAPPING_OUTPUT_BUF) ? (size_t)-1 : ((pOut_buf_next - pOut_buf_start) + *pOut_buf_size) - 1, dist_from_out_buf_start;
 
@@ -959,7 +959,7 @@ int mz_inflateInit2(mz_streamp pStream, int window_bits, inflate_state *pDecomp)
     pStream->reserved = 0;
 
     //pStream->state = (struct mz_internal_state *)pDecomp;
-    pStream->state = (struct inflate_state *) pDecomp;
+    pStream->state = (inflate_state *) pDecomp;
 
     tinfl_init(&pDecomp->m_decomp);
     pDecomp->m_dict_ofs = 0;

OpenCL/m17220_a3-pure.cl

@@ -87,6 +87,7 @@ Related publication: https://scitepress.org/PublicationsDetail.aspx?ID=KLPzPqStp
 
 #include "inc_vendor.h"
 #include "inc_types.h"
+#include "inc_platform.cl"
 #include "inc_common.cl"
 #include "inc_simd.cl"
 
@@ -129,7 +130,7 @@ struct pkzip_hash
   u32 data_length;
   u16 checksum_from_crc;
   u16 checksum_from_timestamp;
-  u8  data[MAX_DATA];
+  u32 data[MAX_DATA];
 
 } __attribute__((packed));
 
@@ -227,7 +228,7 @@ typedef struct {
     u16 val;         /* offset in table or code value */
 } code;
 
-CONSTANT_AS code lenfix[512] = {
+CONSTANT_VK code lenfix[512] = {
     {96,7,0},{0,8,80},{0,8,16},{20,8,115},{18,7,31},{0,8,112},{0,8,48},
     {0,9,192},{16,7,10},{0,8,96},{0,8,32},{0,9,160},{0,8,0},{0,8,128},
     {0,8,64},{0,9,224},{16,7,6},{0,8,88},{0,8,24},{0,9,144},{19,7,59},
@@ -304,7 +305,7 @@ CONSTANT_AS code lenfix[512] = {
     {0,9,255}
 };
 
-CONSTANT_AS code distfix[32] = {
+CONSTANT_VK code distfix[32] = {
     {16,5,1},{23,5,257},{19,5,17},{27,5,4097},{17,5,5},{25,5,1025},
     {21,5,65},{29,5,16385},{16,5,3},{24,5,513},{20,5,33},{28,5,8193},
     {18,5,9},{26,5,2049},{22,5,129},{64,5,0},{16,5,2},{23,5,385},
@@ -538,14 +539,14 @@ __kernel void m17220_sxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
    * sbox, kbox
    */
 
-  LOCAL_AS u32 l_crc32tab[256];
+  LOCAL_VK u32 l_crc32tab[256];
 
   for (u64 i = lid; i < 256; i += lsz)
   {
     l_crc32tab[i] = crc32tab[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS();
 
   if (gid >= gid_max) return;
 
@@ -604,13 +605,11 @@ __kernel void m17220_sxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
 
     for (u32 idx = 0; idx < hash_count; idx++)
     {
-      __global u32 *data_ptr = (__global u32 *) esalt_bufs[digests_offset].hashes[idx].data;
-
       u32x key0 = key0init;
       u32x key1 = key1init;
       u32x key2 = key2init;
 
-      next = data_ptr[0];
+      next = esalt_bufs[digests_offset].hashes[idx].data[0];
 
       update_key3 (key2, key3);
       plain = unpack_v8a_from_v32_S (next) ^ key3;
@@ -628,7 +627,7 @@ __kernel void m17220_sxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
       plain = unpack_v8d_from_v32_S (next) ^ key3;
       update_key012 (key0, key1, key2, plain, l_crc32tab);
 
-      next = data_ptr[1];
+      next = esalt_bufs[digests_offset].hashes[idx].data[1];
 
       update_key3 (key2, key3);
       plain = unpack_v8a_from_v32_S (next) ^ key3;
@@ -646,7 +645,7 @@ __kernel void m17220_sxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
       plain = unpack_v8d_from_v32_S (next) ^ key3;
       update_key012 (key0, key1, key2, plain, l_crc32tab);
 
-      next = data_ptr[2];
+      next = esalt_bufs[digests_offset].hashes[idx].data[2];
 
       update_key3 (key2, key3);
       plain = unpack_v8a_from_v32_S (next) ^ key3;
@@ -672,7 +671,7 @@ __kernel void m17220_sxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
 
       u8 tmp[TMPSIZ];
 
-      next = data_ptr[3];
+      next = esalt_bufs[digests_offset].hashes[idx].data[3];
 
       update_key3 (key2, key3);
       plain = unpack_v8a_from_v32_S (next) ^ key3;
@@ -697,7 +696,7 @@ __kernel void m17220_sxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
 
       for (int i = 16; i < 36; i += 4)
       {
-        next = data_ptr[i / 4];
+        next = esalt_bufs[digests_offset].hashes[idx].data[i / 4];
 
         update_key3 (key2, key3);
         plain = unpack_v8a_from_v32_S (next) ^ key3;
@@ -734,7 +733,7 @@ __kernel void m17220_sxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
 
       infstream.opaque    = Z_NULL;
       infstream.avail_in  = esalt_bufs[digests_offset].hashes[idx].data_length - 12; // size of input
-      infstream.next_in   = esalt_bufs[digests_offset].hashes[idx].data        + 12; // input char array
+      infstream.next_in   = (GLOBAL_AS u8 *) esalt_bufs[digests_offset].hashes[idx].data        + 12; // input char array
       infstream.avail_out = TMPSIZ; // size of output
       infstream.next_out  = tmp; // output char array
 
@@ -805,14 +804,14 @@ __kernel void m17220_mxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
    * sbox, kbox
    */
 
-  LOCAL_AS u32 l_crc32tab[256];
+  LOCAL_VK u32 l_crc32tab[256];
 
   for (u64 i = lid; i < 256; i += lsz)
   {
     l_crc32tab[i] = crc32tab[i];
   }
 
-  barrier (CLK_LOCAL_MEM_FENCE);
+  SYNC_THREADS();
 
   if (gid >= gid_max) return;
 
@@ -871,13 +870,11 @@ __kernel void m17220_mxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
 
     for (u32 idx = 0; idx < hash_count; idx++)
     {
-      __global u32 *data_ptr = (__global u32 *) esalt_bufs[digests_offset].hashes[idx].data;
-
       u32x key0 = key0init;
       u32x key1 = key1init;
       u32x key2 = key2init;
 
-      next = data_ptr[0];
+      next = esalt_bufs[digests_offset].hashes[idx].data[0];
 
       update_key3 (key2, key3);
       plain = unpack_v8a_from_v32_S (next) ^ key3;
@@ -895,7 +892,7 @@ __kernel void m17220_mxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
       plain = unpack_v8d_from_v32_S (next) ^ key3;
       update_key012 (key0, key1, key2, plain, l_crc32tab);
 
-      next = data_ptr[1];
+      next = esalt_bufs[digests_offset].hashes[idx].data[1];
 
       update_key3 (key2, key3);
       plain = unpack_v8a_from_v32_S (next) ^ key3;
@@ -913,7 +910,7 @@ __kernel void m17220_mxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
       plain = unpack_v8d_from_v32_S (next) ^ key3;
       update_key012 (key0, key1, key2, plain, l_crc32tab);
 
-      next = data_ptr[2];
+      next = esalt_bufs[digests_offset].hashes[idx].data[2];
 
       update_key3 (key2, key3);
       plain = unpack_v8a_from_v32_S (next) ^ key3;
@@ -939,7 +936,7 @@ __kernel void m17220_mxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
 
       u8 tmp[TMPSIZ];
 
-      next = data_ptr[3];
+      next = esalt_bufs[digests_offset].hashes[idx].data[3];
 
       update_key3 (key2, key3);
       plain = unpack_v8a_from_v32_S (next) ^ key3;
@@ -964,7 +961,7 @@ __kernel void m17220_mxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
 
       for (int i = 16; i < 36; i += 4)
       {
-        next = data_ptr[i / 4];
+        next = esalt_bufs[digests_offset].hashes[idx].data[i / 4];
 
         update_key3 (key2, key3);
         plain = unpack_v8a_from_v32_S (next) ^ key3;
@@ -990,7 +987,8 @@ __kernel void m17220_mxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
       if (((tmp[0]) & 6) == 2 && !check_inflate_code1 (tmp, 36)) break;
       if (((tmp[0]) & 6) == 4 && !check_inflate_code2 (tmp))     break;
 
-      if (esalt_bufs[digests_offset].hashes[idx].data_type_enum == 1){
+      if (esalt_bufs[digests_offset].hashes[idx].data_type_enum == 1)
+      {
         continue; // so far everything matches for this hash, but it's only a partial one, so we need to continue with the next one
       }
 
@@ -1000,7 +998,7 @@ __kernel void m17220_mxx (KERN_ATTR_VECTOR_ESALT (pkzip_t))
 
       infstream.opaque    = Z_NULL;
       infstream.avail_in  = esalt_bufs[digests_offset].hashes[idx].data_length - 12; // size of input
-      infstream.next_in   = esalt_bufs[digests_offset].hashes[idx].data        + 12; // input char array
+      infstream.next_in   = (GLOBAL_AS u8 *) esalt_bufs[digests_offset].hashes[idx].data        + 12; // input char array
       infstream.avail_out = TMPSIZ; // size of output
       infstream.next_out  = tmp; // output char array
 

src/modules/module_17220.c

@@ -125,7 +125,7 @@ struct pkzip_hash
   u32 data_length;
   u16 checksum_from_crc;
   u16 checksum_from_timestamp;
-  u8  data[MAX_DATA];
+  u32 data[MAX_DATA];
 
 } __attribute__((packed));