2022-02-04 18:54:00 +00:00
|
|
|
/**
|
|
|
|
* Author......: See docs/credits.txt
|
|
|
|
* License.....: MIT
|
|
|
|
*/
|
|
|
|
|
2023-01-17 19:25:40 +00:00
|
|
|
#ifndef INC_MARKOV_H
|
|
|
|
#define INC_MARKOV_H
|
2022-02-04 18:54:00 +00:00
|
|
|
|
|
|
|
#ifdef IS_METAL
|
|
|
|
|
|
|
|
#define KERN_ATTR_L_MARKOV \
|
|
|
|
GLOBAL_AS pw_t *pws_buf_l, \
|
|
|
|
GLOBAL_AS const cs_t *root_css_buf, \
|
|
|
|
GLOBAL_AS const cs_t *markov_css_buf, \
|
|
|
|
CONSTANT_AS const u64 &off, \
|
|
|
|
CONSTANT_AS const u32 &pw_l_len, \
|
|
|
|
CONSTANT_AS const u32 &pw_r_len, \
|
|
|
|
CONSTANT_AS const u32 &mask80, \
|
|
|
|
CONSTANT_AS const u32 &bits14, \
|
|
|
|
CONSTANT_AS const u32 &bits15, \
|
|
|
|
CONSTANT_AS const u64 &gid_max, \
|
|
|
|
uint hc_gid [[ thread_position_in_grid ]]
|
|
|
|
|
|
|
|
#define KERN_ATTR_R_MARKOV \
|
|
|
|
GLOBAL_AS bf_t *pws_buf_r, \
|
|
|
|
GLOBAL_AS const cs_t *root_css_buf, \
|
|
|
|
GLOBAL_AS const cs_t *markov_css_buf, \
|
|
|
|
CONSTANT_AS const u64 &off, \
|
|
|
|
CONSTANT_AS const u32 &pw_r_len, \
|
|
|
|
CONSTANT_AS const u32 &mask80, \
|
|
|
|
CONSTANT_AS const u32 &bits14, \
|
|
|
|
CONSTANT_AS const u32 &bits15, \
|
|
|
|
CONSTANT_AS const u64 &gid_max, \
|
|
|
|
uint hc_gid [[ thread_position_in_grid ]]
|
|
|
|
|
|
|
|
#define KERN_ATTR_C_MARKOV \
|
|
|
|
GLOBAL_AS pw_t *pws_buf, \
|
|
|
|
GLOBAL_AS const cs_t *root_css_buf, \
|
|
|
|
GLOBAL_AS const cs_t *markov_css_buf, \
|
|
|
|
CONSTANT_AS const u64 &off, \
|
|
|
|
CONSTANT_AS const u32 &pw_len, \
|
|
|
|
CONSTANT_AS const u32 &mask80, \
|
|
|
|
CONSTANT_AS const u32 &bits14, \
|
|
|
|
CONSTANT_AS const u32 &bits15, \
|
|
|
|
CONSTANT_AS const u64 &gid_max, \
|
|
|
|
uint hc_gid [[ thread_position_in_grid ]]
|
|
|
|
|
|
|
|
#else // CUDA, HIP, OpenCL
|
|
|
|
|
|
|
|
#define KERN_ATTR_L_MARKOV \
|
|
|
|
GLOBAL_AS pw_t *pws_buf_l, \
|
|
|
|
GLOBAL_AS const cs_t *root_css_buf, \
|
|
|
|
GLOBAL_AS const cs_t *markov_css_buf, \
|
|
|
|
const u64 off, \
|
|
|
|
const u32 pw_l_len, \
|
|
|
|
const u32 pw_r_len, \
|
|
|
|
const u32 mask80, \
|
|
|
|
const u32 bits14, \
|
|
|
|
const u32 bits15, \
|
|
|
|
const u64 gid_max
|
|
|
|
|
|
|
|
#define KERN_ATTR_R_MARKOV \
|
|
|
|
GLOBAL_AS bf_t *pws_buf_r, \
|
|
|
|
GLOBAL_AS const cs_t *root_css_buf, \
|
|
|
|
GLOBAL_AS const cs_t *markov_css_buf, \
|
|
|
|
const u64 off, \
|
|
|
|
const u32 pw_r_len, \
|
|
|
|
const u32 mask80, \
|
|
|
|
const u32 bits14, \
|
|
|
|
const u32 bits15, \
|
|
|
|
const u64 gid_max
|
|
|
|
|
|
|
|
#define KERN_ATTR_C_MARKOV \
|
|
|
|
GLOBAL_AS pw_t *pws_buf, \
|
|
|
|
GLOBAL_AS const cs_t *root_css_buf, \
|
|
|
|
GLOBAL_AS const cs_t *markov_css_buf, \
|
|
|
|
const u64 off, \
|
|
|
|
const u32 pw_len, \
|
|
|
|
const u32 mask80, \
|
|
|
|
const u32 bits14, \
|
|
|
|
const u32 bits15, \
|
|
|
|
const u64 gid_max
|
|
|
|
|
|
|
|
#endif // IS_METAL
|
|
|
|
|
2023-01-17 19:25:40 +00:00
|
|
|
#endif // INC_MARKOV_H
|
2022-02-04 18:54:00 +00:00
|
|
|
|