From c9c7261a050fb7711c7c0d3e46073b9505342c57 Mon Sep 17 00:00:00 2001 From: Sein Coray Date: Tue, 14 May 2019 14:38:55 +0200 Subject: [PATCH] updated m17210 to be cuda compatible --- OpenCL/m17210_a0-pure.cl | 31 ++++++++++++++----------------- OpenCL/m17210_a1-pure.cl | 31 ++++++++++++++----------------- OpenCL/m17210_a3-pure.cl | 31 ++++++++++++++----------------- 3 files changed, 42 insertions(+), 51 deletions(-) diff --git a/OpenCL/m17210_a0-pure.cl b/OpenCL/m17210_a0-pure.cl index 8308d10ef..e6f4515ad 100644 --- a/OpenCL/m17210_a0-pure.cl +++ b/OpenCL/m17210_a0-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" #include "inc_rp.h" @@ -131,7 +132,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)); @@ -219,7 +220,7 @@ CONSTANT_AS u32a crc32tab[256] = 0xb40bbe37, 0xc30c8ea1, 0x5a05df1b, 0x2d02ef8d }; -__kernel void m17210_sxx (KERN_ATTR_RULES_ESALT (pkzip_t)) +KERNEL_FQ void m17210_sxx (KERN_ATTR_RULES_ESALT (pkzip_t)) { /** * modifier @@ -233,25 +234,23 @@ __kernel void m17210_sxx (KERN_ATTR_RULES_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(); - __global u32 *data_ptr = (__global u32 *) esalt_bufs[digests_offset].hash.data; - - LOCAL_AS u32 l_data[MAX_LOCAL]; + LOCAL_VK u32 l_data[MAX_LOCAL]; for (u64 i = lid; i < MAX_LOCAL; i += lsz) { - l_data[i] = data_ptr[i]; + l_data[i] = esalt_bufs[digests_offset].hash.data[i]; } - barrier (CLK_LOCAL_MEM_FENCE); + SYNC_THREADS(); if (gid >= gid_max) return; @@ -428,7 +427,7 @@ __kernel void m17210_sxx (KERN_ATTR_RULES_ESALT (pkzip_t)) } } -__kernel void m17210_mxx (KERN_ATTR_RULES_ESALT (pkzip_t)) +KERNEL_FQ void m17210_mxx (KERN_ATTR_RULES_ESALT (pkzip_t)) { /** * modifier @@ -442,25 +441,23 @@ __kernel void m17210_mxx (KERN_ATTR_RULES_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); - - __global u32 *data_ptr = (__global u32 *) esalt_bufs[digests_offset].hash.data; + SYNC_THREADS(); - LOCAL_AS u32 l_data[MAX_LOCAL]; + LOCAL_VK u32 l_data[MAX_LOCAL]; for (u64 i = lid; i < MAX_LOCAL; i += lsz) { - l_data[i] = data_ptr[i]; + l_data[i] = esalt_bufs[digests_offset].hash.data[i]; } - barrier (CLK_LOCAL_MEM_FENCE); + SYNC_THREADS(); if (gid >= gid_max) return; diff --git a/OpenCL/m17210_a1-pure.cl b/OpenCL/m17210_a1-pure.cl index 4e726f759..4540b20b2 100644 --- a/OpenCL/m17210_a1-pure.cl +++ b/OpenCL/m17210_a1-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)); @@ -217,7 +218,7 @@ CONSTANT_AS u32a crc32tab[256] = 0xb40bbe37, 0xc30c8ea1, 0x5a05df1b, 0x2d02ef8d }; -__kernel void m17210_sxx (KERN_ATTR_ESALT (pkzip_t)) +KERNEL_FQ void m17210_sxx (KERN_ATTR_ESALT (pkzip_t)) { /** * modifier @@ -231,25 +232,23 @@ __kernel void m17210_sxx (KERN_ATTR_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(); - __global u32 *data_ptr = (__global u32 *) esalt_bufs[digests_offset].hash.data; - - LOCAL_AS u32 l_data[MAX_LOCAL]; + LOCAL_VK u32 l_data[MAX_LOCAL]; for (u64 i = lid; i < MAX_LOCAL; i += lsz) { - l_data[i] = data_ptr[i]; + l_data[i] = esalt_bufs[digests_offset].hash.data[i]; } - barrier (CLK_LOCAL_MEM_FENCE); + SYNC_THREADS(); if (gid >= gid_max) return; @@ -428,7 +427,7 @@ __kernel void m17210_sxx (KERN_ATTR_ESALT (pkzip_t)) } } -__kernel void m17210_mxx (KERN_ATTR_ESALT (pkzip_t)) +KERNEL_FQ void m17210_mxx (KERN_ATTR_ESALT (pkzip_t)) { /** * modifier @@ -442,25 +441,23 @@ __kernel void m17210_mxx (KERN_ATTR_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); - - __global u32 *data_ptr = (__global u32 *) esalt_bufs[digests_offset].hash.data; + SYNC_THREADS(); - LOCAL_AS u32 l_data[MAX_LOCAL]; + LOCAL_VK u32 l_data[MAX_LOCAL]; for (u64 i = lid; i < MAX_LOCAL; i += lsz) { - l_data[i] = data_ptr[i]; + l_data[i] = esalt_bufs[digests_offset].hash.data[i]; } - barrier (CLK_LOCAL_MEM_FENCE); + SYNC_THREADS(); if (gid >= gid_max) return; diff --git a/OpenCL/m17210_a3-pure.cl b/OpenCL/m17210_a3-pure.cl index c15fb57db..e5796ecab 100644 --- a/OpenCL/m17210_a3-pure.cl +++ b/OpenCL/m17210_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)); @@ -217,7 +218,7 @@ CONSTANT_AS u32a crc32tab[256] = 0xb40bbe37, 0xc30c8ea1, 0x5a05df1b, 0x2d02ef8d }; -__kernel void m17210_sxx (KERN_ATTR_VECTOR_ESALT (pkzip_t)) +KERNEL_FQ void m17210_sxx (KERN_ATTR_VECTOR_ESALT (pkzip_t)) { /** * modifier @@ -231,25 +232,23 @@ __kernel void m17210_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(); - __global u32 *data_ptr = (__global u32 *) esalt_bufs[digests_offset].hash.data; - - LOCAL_AS u32 l_data[MAX_LOCAL]; + LOCAL_VK u32 l_data[MAX_LOCAL]; for (u64 i = lid; i < MAX_LOCAL; i += lsz) { - l_data[i] = data_ptr[i]; + l_data[i] = esalt_bufs[digests_offset].hash.data[i]; } - barrier (CLK_LOCAL_MEM_FENCE); + SYNC_THREADS(); if (gid >= gid_max) return; @@ -440,7 +439,7 @@ __kernel void m17210_sxx (KERN_ATTR_VECTOR_ESALT (pkzip_t)) } } -__kernel void m17210_mxx (KERN_ATTR_VECTOR_ESALT (pkzip_t)) +KERNEL_FQ void m17210_mxx (KERN_ATTR_VECTOR_ESALT (pkzip_t)) { /** * modifier @@ -454,25 +453,23 @@ __kernel void m17210_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); - - __global u32 *data_ptr = (__global u32 *) esalt_bufs[digests_offset].hash.data; + SYNC_THREADS(); - LOCAL_AS u32 l_data[MAX_LOCAL]; + LOCAL_VK u32 l_data[MAX_LOCAL]; for (u64 i = lid; i < MAX_LOCAL; i += lsz) { - l_data[i] = data_ptr[i]; + l_data[i] = esalt_bufs[digests_offset].hash.data[i]; } - barrier (CLK_LOCAL_MEM_FENCE); + SYNC_THREADS(); if (gid >= gid_max) return;