GPU рдХреЗ рд▓рд┐рдП рд╕рд░рд▓ рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛

GPU рдХреЗ рд▓рд┐рдП рд╕рд░рд▓ рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛
рдореИрдВрдиреЗ рдЗрд╕реЗ рдЬреАрдердм рдкрд░ рдкреЛрд╕реНрдЯ рдХрд┐рдпрд╛ рдирдпрд╛ рдкреНрд░реЛрдЬреЗрдХреНрдЯ рдПрдХ рд╕рд░рд▓ рдЬреАрдкреАрдпреВ рд╣реИрд╢ рдЯреЗрдмрд▓.

рдпрд╣ рдПрдХ рд╕рд╛рдзрд╛рд░рдг рдЬреАрдкреАрдпреВ рд╣реИрд╢ рдЯреЗрдмрд▓ рд╣реИ рдЬреЛ рдкреНрд░рддрд┐ рд╕реЗрдХрдВрдб рд╕реИрдХрдбрд╝реЛрдВ рд▓рд╛рдЦреЛрдВ рдЗрдВрд╕рд░реНрдЯ рдХреЛ рдкреНрд░реЛрд╕реЗрд╕ рдХрд░рдиреЗ рдореЗрдВ рд╕рдХреНрд╖рдо рд╣реИред рдореЗрд░реЗ NVIDIA GTX 1060 рд▓реИрдкрдЯреЙрдк рдкрд░, рдХреЛрдб рд▓рдЧрднрдЧ 64 рдПрдордПрд╕ рдореЗрдВ 210 рдорд┐рд▓рд┐рдпрди рдмреЗрддрд░рддреАрдм рдврдВрдЧ рд╕реЗ рдЙрддреНрдкрдиреНрди рдХреБрдВрдЬреА-рдореВрд▓реНрдп рдЬреЛрдбрд╝реЗ рдбрд╛рд▓рддрд╛ рд╣реИ рдФрд░ рд▓рдЧрднрдЧ 32 рдПрдордПрд╕ рдореЗрдВ 64 рдорд┐рд▓рд┐рдпрди рдЬреЛрдбрд╝реЗ рд╣рдЯрд╛ рджреЗрддрд╛ рд╣реИред

рдпрд╛рдиреА, рд▓реИрдкрдЯреЙрдк рдкрд░ рд╕реНрдкреАрдб рд▓рдЧрднрдЧ 300 рдорд┐рд▓рд┐рдпрди рдЗрдВрд╕рд░реНрдЯ/рд╕реЗрдХрдВрдб рдФрд░ 500 рдорд┐рд▓рд┐рдпрди рдбрд┐рд▓реАрдЯ/рд╕реЗрдХрдВрдб рд╣реИред

рддрд╛рд▓рд┐рдХрд╛ CUDA рдореЗрдВ рд▓рд┐рдЦреА рдЧрдИ рд╣реИ, рд╣рд╛рд▓рд╛рдБрдХрд┐ рд╡рд╣реА рддрдХрдиреАрдХ HLSL рдпрд╛ GLSL рдкрд░ рд▓рд╛рдЧреВ рдХреА рдЬрд╛ рд╕рдХрддреА рд╣реИред рд╡реАрдбрд┐рдпреЛ рдХрд╛рд░реНрдб рдкрд░ рдЙрдЪреНрдЪ рдкреНрд░рджрд░реНрд╢рди рд╕реБрдирд┐рд╢реНрдЪрд┐рдд рдХрд░рдиреЗ рдХреЗ рд▓рд┐рдП рдХрд╛рд░реНрдпрд╛рдиреНрд╡рдпрди рдХреА рдХрдИ рд╕реАрдорд╛рдПрдБ рд╣реИрдВ:

  • рдХреЗрд╡рд▓ 32-рдмрд┐рдЯ рдХреБрдВрдЬрд┐рдпрд╛рдБ рдФрд░ рд╕рдорд╛рди рдорд╛рди рд╕рдВрд╕рд╛рдзрд┐рдд рд╣реЛрддреЗ рд╣реИрдВред
  • рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдХрд╛ рдПрдХ рдирд┐рд╢реНрдЪрд┐рдд рдЖрдХрд╛рд░ рд╣реЛрддрд╛ рд╣реИред
  • рдФрд░ рдпрд╣ рдЖрдХрд╛рд░ рджреЛ рдШрд╛рдд рдХреЗ рдмрд░рд╛рдмрд░ рд╣реЛрдирд╛ рдЪрд╛рд╣рд┐рдПред

рдХреБрдВрдЬрд┐рдпреЛрдВ рдФрд░ рдорд╛рдиреЛрдВ рдХреЗ рд▓рд┐рдП, рдЖрдкрдХреЛ рдПрдХ рд╕рд╛рдзрд╛рд░рдг рд╕реАрдорд╛рдВрдХрдХ рдорд╛рд░реНрдХрд░ рдЖрд░рдХреНрд╖рд┐рдд рдХрд░рдирд╛ рд╣реЛрдЧрд╛ (рдЙрдкрд░реЛрдХреНрдд рдХреЛрдб рдореЗрдВ рдпрд╣ 0xffffffff рд╣реИ)ред

рддрд╛рд▓реЗ рдХреЗ рдмрд┐рдирд╛ рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛

рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдУрдкрди рдПрдбреНрд░реЗрд╕рд┐рдВрдЧ рдХрд╛ рдЙрдкрдпреЛрдЧ рдХрд░рддреА рд╣реИ рд░реИрдЦрд┐рдХ рдЬрд╛рдВрдЪ, рдпрд╛рдиреА, рдпрд╣ рдХреЗрд╡рд▓ рдХреБрдВрдЬреА-рдореВрд▓реНрдп рдЬреЛрдбрд╝реЗ рдХреА рдПрдХ рд╕рд░рдгреА рд╣реИ рдЬреЛ рдореЗрдореЛрд░реА рдореЗрдВ рд╕рдВрдЧреНрд░рд╣реАрдд рд╣реЛрддреА рд╣реИ рдФрд░ рдЗрд╕рдореЗрдВ рдмреЗрд╣рддрд░ рдХреИрд╢ рдкреНрд░рджрд░реНрд╢рди рд╣реЛрддрд╛ рд╣реИред рдЪреЗрдирд┐рдВрдЧ рдХреЗ рд▓рд┐рдП рдРрд╕рд╛ рдирд╣реАрдВ рдХрд╣рд╛ рдЬрд╛ рд╕рдХрддрд╛ рд╣реИ, рдЬрд┐рд╕рдореЗрдВ рд▓рд┐рдВрдХ рдХреА рдЧрдИ рд╕реВрдЪреА рдореЗрдВ рдкреЙрдЗрдВрдЯрд░ рдХреА рдЦреЛрдЬ рдХрд░рдирд╛ рд╢рд╛рдорд┐рд▓ рд╣реИред рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рддрддреНрд╡реЛрдВ рдХреЛ рд╕рдВрдЧреНрд░рд╣рд┐рдд рдХрд░рдиреЗ рд╡рд╛рд▓реА рдПрдХ рд╕рд░рд▓ рд╕рд╛рд░рдгреА рд╣реИ KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

рддрд╛рд▓рд┐рдХрд╛ рдХрд╛ рдЖрдХрд╛рд░ рджреЛ рдХреА рдШрд╛рдд рд╣реИ, рдЕрднрд╛рдЬреНрдп рд╕рдВрдЦреНрдпрд╛ рдирд╣реАрдВ, рдХреНрдпреЛрдВрдХрд┐ рдПрдХ рддреЗрдЬрд╝ рдирд┐рд░реНрджреЗрд╢ pow2/AND рдорд╛рд╕реНрдХ рд▓рдЧрд╛рдиреЗ рдХреЗ рд▓рд┐рдП рдкрд░реНрдпрд╛рдкреНрдд рд╣реИ, рд▓реЗрдХрд┐рди рдорд╛рдкрд╛рдВрдХ рдСрдкрд░реЗрдЯрд░ рдмрд╣реБрдд рдзреАрдорд╛ рд╣реИред рд░реИрдЦрд┐рдХ рдЬрд╛рдВрдЪ рдХреЗ рдорд╛рдорд▓реЗ рдореЗрдВ рдпрд╣ рдорд╣рддреНрд╡рдкреВрд░реНрдг рд╣реИ, рдХреНрдпреЛрдВрдХрд┐ рдПрдХ рд░реИрдЦрд┐рдХ рддрд╛рд▓рд┐рдХрд╛ рд▓реБрдХрдЕрдк рдореЗрдВ рд╕реНрд▓реЙрдЯ рдЗрдВрдбреЗрдХреНрд╕ рдХреЛ рдкреНрд░рддреНрдпреЗрдХ рд╕реНрд▓реЙрдЯ рдореЗрдВ рд▓рдкреЗрдЯрд╛ рдЬрд╛рдирд╛ рдЪрд╛рд╣рд┐рдПред рдФрд░ рдкрд░рд┐рдгрд╛рдорд╕реНрд╡рд░реВрдк, рдкреНрд░рддреНрдпреЗрдХ рд╕реНрд▓реЙрдЯ рдореЗрдВ рдСрдкрд░реЗрд╢рди рдХреА рд▓рд╛рдЧрдд рдореЙрдбреНрдпреВрд▓реЛ рдореЗрдВ рдЬреЛрдбрд╝ рджреА рдЬрд╛рддреА рд╣реИред

рддрд╛рд▓рд┐рдХрд╛ рдХреЗрд╡рд▓ рдкреНрд░рддреНрдпреЗрдХ рддрддреНрд╡ рдХреЗ рд▓рд┐рдП рдХреБрдВрдЬреА рдФрд░ рдорд╛рди рд╕рдВрдЧреНрд░рд╣реАрдд рдХрд░рддреА рд╣реИ, рдХреБрдВрдЬреА рдХрд╛ рд╣реИрд╢ рдирд╣реАрдВред рдЪреВрдБрдХрд┐ рддрд╛рд▓рд┐рдХрд╛ рдХреЗрд╡рд▓ 32-рдмрд┐рдЯ рдХреБрдВрдЬрд┐рдпрд╛рдБ рд╕рдВрдЧреНрд░рд╣реАрдд рдХрд░рддреА рд╣реИ, рдЗрд╕рд▓рд┐рдП рд╣реИрд╢ рдХреА рдЧрдгрдирд╛ рдмрд╣реБрдд рдЬрд▓реНрджреА рдХреА рдЬрд╛рддреА рд╣реИред рдЙрдкрд░реЛрдХреНрдд рдХреЛрдб Murmur3 рд╣реИрд╢ рдХрд╛ рдЙрдкрдпреЛрдЧ рдХрд░рддрд╛ рд╣реИ, рдЬреЛ рдХреЗрд╡рд▓ рдХреБрдЫ рдмрджрд▓рд╛рд╡, XOR рдФрд░ рдЧреБрдгрди рдХрд░рддрд╛ рд╣реИред

рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рд▓реЙрдХрд┐рдВрдЧ рд╕реБрд░рдХреНрд╖рд╛ рддрдХрдиреАрдХреЛрдВ рдХрд╛ рдЙрдкрдпреЛрдЧ рдХрд░рддреА рд╣реИ рдЬреЛ рдореЗрдореЛрд░реА рдСрд░реНрдбрд░ рд╕реЗ рд╕реНрд╡рддрдВрддреНрд░ рд╣реЛрддреА рд╣реИрдВред рднрд▓реЗ рд╣реА рдХреБрдЫ рд▓реЗрдЦрди рдСрдкрд░реЗрд╢рди рдРрд╕реЗ рдЕрдиреНрдп рдСрдкрд░реЗрд╢рдиреЛрдВ рдХреЗ рдХреНрд░рдо рдХреЛ рдмрд╛рдзрд┐рдд рдХрд░рддреЗ рд╣реЛрдВ, рдлрд┐рд░ рднреА рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рд╕рд╣реА рд╕реНрдерд┐рддрд┐ рдмрдирд╛рдП рд░рдЦреЗрдЧреАред рд╣рдо рдЗрд╕рдХреЗ рдмрд╛рд░реЗ рдореЗрдВ рдиреАрдЪреЗ рдмрд╛рдд рдХрд░реЗрдВрдЧреЗ. рдпрд╣ рддрдХрдиреАрдХ рдЙрди рд╡реАрдбрд┐рдпреЛ рдХрд╛рд░реНрдбреЛрдВ рдХреЗ рд╕рд╛рде рдмрдврд╝рд┐рдпрд╛ рдХрд╛рдо рдХрд░рддреА рд╣реИ рдЬреЛ рд╣рдЬрд╛рд░реЛрдВ рдереНрд░реЗрдб рдПрдХ рд╕рд╛рде рдЪрд▓рд╛рддреЗ рд╣реИрдВред

рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдореЗрдВ рдХреБрдВрдЬрд┐рдпрд╛рдБ рдФрд░ рдорд╛рди рдЦрд╛рд▓реА рдХрд░рдиреЗ рдХреЗ рд▓рд┐рдП рдкреНрд░рд╛рд░рдВрдн рдХрд┐рдП рдЬрд╛рддреЗ рд╣реИрдВред

рдХреЛрдб рдХреЛ 64-рдмрд┐рдЯ рдХреБрдВрдЬрд┐рдпреЛрдВ рдФрд░ рдорд╛рдиреЛрдВ рдХреЛ рд╕рдВрднрд╛рд▓рдиреЗ рдХреЗ рд▓рд┐рдП рднреА рд╕рдВрд╢реЛрдзрд┐рдд рдХрд┐рдпрд╛ рдЬрд╛ рд╕рдХрддрд╛ рд╣реИред рдХреБрдВрдЬрд┐рдпреЛрдВ рдХреЛ рдкрд░рдорд╛рдгреБ рдкрдврд╝рдиреЗ, рд▓рд┐рдЦрдиреЗ рдФрд░ рддреБрд▓рдирд╛-рдФрд░-рд╕реНрд╡реИрдк рд╕рдВрдЪрд╛рд▓рди рдХреА рдЖрд╡рд╢реНрдпрдХрддрд╛ рд╣реЛрддреА рд╣реИред рдФрд░ рдореВрд▓реНрдпреЛрдВ рдХреЗ рд▓рд┐рдП рдкрд░рдорд╛рдгреБ рдкрдврд╝рдиреЗ рдФрд░ рд▓рд┐рдЦрдиреЗ рдХреЗ рд╕рдВрдЪрд╛рд▓рди рдХреА рдЖрд╡рд╢реНрдпрдХрддрд╛ рд╣реЛрддреА рд╣реИред рд╕реМрднрд╛рдЧреНрдп рд╕реЗ, CUDA рдореЗрдВ, 32- рдФрд░ 64-рдмрд┐рдЯ рдорд╛рдиреЛрдВ рдХреЗ рд▓рд┐рдП рдкрдврд╝рдиреЗ-рд▓рд┐рдЦрдиреЗ рдХреЗ рд╕рдВрдЪрд╛рд▓рди рддрдм рддрдХ рдкрд░рдорд╛рдгреБ рд╣реЛрддреЗ рд╣реИрдВ рдЬрдм рддрдХ рд╡реЗ рд╕реНрд╡рд╛рднрд╛рд╡рд┐рдХ рд░реВрдк рд╕реЗ рд╕рдВрд░реЗрдЦрд┐рдд рд╣реЛрддреЗ рд╣реИрдВ (рдиреАрдЪреЗ рджреЗрдЦреЗрдВ)ред рдпрд╣рд╛рдВ), рдФрд░ рдЖрдзреБрдирд┐рдХ рд╡реАрдбрд┐рдпреЛ рдХрд╛рд░реНрдб 64-рдмрд┐рдЯ рдкрд░рдорд╛рдгреБ рддреБрд▓рдирд╛-рдФрд░-рд╡рд┐рдирд┐рдордп рд╕рдВрдЪрд╛рд▓рди рдХрд╛ рд╕рдорд░реНрдерди рдХрд░рддреЗ рд╣реИрдВред рдмреЗрд╢рдХ, 64 рдмрд┐рдЯреНрд╕ рдкрд░ рдЬрд╛рдиреЗ рдкрд░ рдкреНрд░рджрд░реНрд╢рди рдереЛрдбрд╝рд╛ рдХрдо рд╣реЛ рдЬрд╛рдПрдЧрд╛ред

рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рд╕реНрдерд┐рддрд┐

рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдореЗрдВ рдкреНрд░рддреНрдпреЗрдХ рдХреБрдВрдЬреА-рдореВрд▓реНрдп рдЬреЛрдбрд╝реА рдореЗрдВ рдЪрд╛рд░ рдЕрд╡рд╕реНрдерд╛рдУрдВ рдореЗрдВ рд╕реЗ рдПрдХ рд╣реЛ рд╕рдХрддреА рд╣реИ:

  • рдХреБрдВрдЬреА рдФрд░ рдорд╛рди рд░рд┐рдХреНрдд рд╣реИрдВ. рдЗрд╕ рд╕реНрдерд┐рддрд┐ рдореЗрдВ, рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдкреНрд░рд╛рд░рдВрдн рдХреА рдЬрд╛рддреА рд╣реИред
  • рдХреБрдВрдЬреА рддреЛ рд▓рд┐рдЦ рджреА рдЧрдИ рд╣реИ, рдкрд░рдВрддреБ рдореВрд▓реНрдп рдЕрднреА рддрдХ рдирд╣реАрдВ рд▓рд┐рдЦрд╛ рдЧрдпрд╛ рд╣реИред рдпрджрд┐ рдХреЛрдИ рдЕрдиреНрдп рдереНрд░реЗрдб рд╡рд░реНрддрдорд╛рди рдореЗрдВ рдбреЗрдЯрд╛ рдкрдврд╝ рд░рд╣рд╛ рд╣реИ, рддреЛ рд╡рд╣ рдЦрд╛рд▓реА рд▓реМрдЯрддрд╛ рд╣реИред рдпрд╣ рд╕рд╛рдорд╛рдиреНрдп рд╣реИ, рдпрджрд┐ рдирд┐рд╖реНрдкрд╛рджрди рдХрд╛ рдХреЛрдИ рдЕрдиреНрдп рдереНрд░реЗрдб рдереЛрдбрд╝рд╛ рдкрд╣рд▓реЗ рдХрд╛рдо рдХрд░рддрд╛ рддреЛ рдпрд╣реА рдмрд╛рдд рд╣реЛрддреА, рдФрд░ рд╣рдо рд╕рдорд╡рд░реНрддреА рдбреЗрдЯрд╛ рд╕рдВрд░рдЪрдирд╛ рдХреЗ рдмрд╛рд░реЗ рдореЗрдВ рдмрд╛рдд рдХрд░ рд░рд╣реЗ рд╣реИрдВред
  • рдХреБрдВрдЬреА рдФрд░ рдорд╛рди рджреЛрдиреЛрдВ рд░рд┐рдХреЙрд░реНрдб рдХрд┐рдП рдЬрд╛рддреЗ рд╣реИрдВ.
  • рдорд╛рди рдирд┐рд╖реНрдкрд╛рджрди рдХреЗ рдЕрдиреНрдп рдереНрд░реЗрдбреНрд╕ рдХреЗ рд▓рд┐рдП рдЙрдкрд▓рдмреНрдз рд╣реИ, рд▓реЗрдХрд┐рди рдХреБрдВрдЬреА рдЕрднреА рддрдХ рдЙрдкрд▓рдмреНрдз рдирд╣реАрдВ рд╣реИред рдРрд╕рд╛ рдЗрд╕рд▓рд┐рдП рд╣реЛ рд╕рдХрддрд╛ рд╣реИ рдХреНрдпреЛрдВрдХрд┐ CUDA рдкреНрд░реЛрдЧреНрд░рд╛рдорд┐рдВрдЧ рдореЙрдбрд▓ рдореЗрдВ рд╢рд┐рдерд┐рд▓ рд░реВрдк рд╕реЗ рдСрд░реНрдбрд░ рдХрд┐рдпрд╛ рдЧрдпрд╛ рдореЗрдореЛрд░реА рдореЙрдбрд▓ рд╣реИред рдпрд╣ рд╕рд╛рдорд╛рдиреНрдп рд╣реИ; рдХрд┐рд╕реА рднреА рд╕реНрдерд┐рддрд┐ рдореЗрдВ, рдХреБрдВрдЬреА рдЕрднреА рднреА рдЦрд╛рд▓реА рд╣реИ, рднрд▓реЗ рд╣реА рдорд╛рди рдЕрдм рдРрд╕рд╛ рди рд╣реЛред

рдПрдХ рдорд╣рддреНрд╡рдкреВрд░реНрдг рдмрд╛рд░реАрдХрд┐рдпрд╛рдВ рдпрд╣ рд╣реИ рдХрд┐ рдПрдХ рдмрд╛рд░ рдХреБрдВрдЬреА рдХреЛ рд╕реНрд▓реЙрдЯ рдореЗрдВ рд▓рд┐рдЦ рджрд┐рдпрд╛ рдЧрдпрд╛ рд╣реИ, рддреЛ рдпрд╣ рдЖрдЧреЗ рдирд╣реАрдВ рдмрдврд╝рддреА рд╣реИ - рднрд▓реЗ рд╣реА рдХреБрдВрдЬреА рд╣рдЯрд╛ рджреА рдЧрдИ рд╣реЛ, рд╣рдо рдЗрд╕рдХреЗ рдмрд╛рд░реЗ рдореЗрдВ рдиреАрдЪреЗ рдмрд╛рдд рдХрд░реЗрдВрдЧреЗред

рд╣реИрд╢ рдЯреЗрдмрд▓ рдХреЛрдб рд╢рд┐рдерд┐рд▓ рдХреНрд░рдо рд╡рд╛рд▓реЗ рдореЗрдореЛрд░реА рдореЙрдбрд▓ рдХреЗ рд╕рд╛рде рднреА рдХрд╛рдо рдХрд░рддрд╛ рд╣реИ рдЬрд┐рд╕рдореЗрдВ рдореЗрдореЛрд░реА рдХреЛ рдкрдврд╝рдиреЗ рдФрд░ рд▓рд┐рдЦрдиреЗ рдХрд╛ рдХреНрд░рдо рдЕрдЬреНрдЮрд╛рдд рд╣реИред рдЬреИрд╕рд╛ рдХрд┐ рд╣рдо рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдореЗрдВ рд╕рдореНрдорд┐рд▓рди, рд▓реБрдХрдЕрдк рдФрд░ рд╡рд┐рд▓реЛрдкрди рдХреЛ рджреЗрдЦрддреЗ рд╣реИрдВ, рдпрд╛рдж рд░рдЦреЗрдВ рдХрд┐ рдкреНрд░рддреНрдпреЗрдХ рдХреБрдВрдЬреА-рдореВрд▓реНрдп рдЬреЛрдбрд╝реА рдКрдкрд░ рд╡рд░реНрдгрд┐рдд рдЪрд╛рд░ рд╕реНрдерд┐рддрд┐рдпреЛрдВ рдореЗрдВ рд╕реЗ рдПрдХ рдореЗрдВ рд╣реИред

рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдореЗрдВ рд╕рдореНрдорд┐рд▓рд┐рдд рдХрд░рдирд╛

CUDA рдлрд╝рдВрдХреНрд╢рди рдЬреЛ рдХреБрдВрдЬреА-рдореВрд▓реНрдп рдЬреЛрдбрд╝реЗ рдХреЛ рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдореЗрдВ рд╕рдореНрдорд┐рд▓рд┐рдд рдХрд░рддрд╛ рд╣реИ, рдЗрд╕ рддрд░рд╣ рджрд┐рдЦрддрд╛ рд╣реИ:

void gpu_hashtable_insert(KeyValue* hashtable, uint32_t key, uint32_t value)
{
    uint32_t slot = hash(key);

    while (true)
    {
        uint32_t prev = atomicCAS(&hashtable[slot].key, kEmpty, key);
        if (prev == kEmpty || prev == key)
        {
            hashtable[slot].value = value;
            break;
        }
        slot = (slot + 1) & (kHashTableCapacity-1);
    }
}

рдПрдХ рдХреБрдВрдЬреА рдбрд╛рд▓рдиреЗ рдХреЗ рд▓рд┐рдП, рдХреЛрдб рд╕рдореНрдорд┐рд▓рд┐рдд рдХреБрдВрдЬреА рдХреЗ рд╣реИрд╢ рд╕реЗ рд╢реБрд░реВ рд╣реЛрдХрд░ рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рд╕рд░рдгреА рдХреЗ рдорд╛рдзреНрдпрдо рд╕реЗ рдкреБрдирд░рд╛рд╡реГрддреНрдд рд╣реЛрддрд╛ рд╣реИред рд╕рд░рдгреА рдореЗрдВ рдкреНрд░рддреНрдпреЗрдХ рд╕реНрд▓реЙрдЯ рдПрдХ рдкрд░рдорд╛рдгреБ рддреБрд▓рдирд╛-рдФрд░-рд╕реНрд╡реИрдк рдСрдкрд░реЗрд╢рди рдХрд░рддрд╛ рд╣реИ рдЬреЛ рдЙрд╕ рд╕реНрд▓реЙрдЯ рдореЗрдВ рдХреБрдВрдЬреА рдХреА рддреБрд▓рдирд╛ рдЦрд╛рд▓реА рд╕реЗ рдХрд░рддрд╛ рд╣реИред рдпрджрд┐ рдмреЗрдореЗрд▓ рдХрд╛ рдкрддрд╛ рдЪрд▓рддрд╛ рд╣реИ, рддреЛ рд╕реНрд▓реЙрдЯ рдореЗрдВ рдХреБрдВрдЬреА рдХреЛ рд╕рдореНрдорд┐рд▓рд┐рдд рдХреБрдВрдЬреА рдХреЗ рд╕рд╛рде рдЕрджреНрдпрддрди рдХрд┐рдпрд╛ рдЬрд╛рддрд╛ рд╣реИ, рдФрд░ рдлрд┐рд░ рдореВрд▓ рд╕реНрд▓реЙрдЯ рдХреБрдВрдЬреА рд╡рд╛рдкрд╕ рдХрд░ рджреА рдЬрд╛рддреА рд╣реИред рдпрджрд┐ рдпрд╣ рдореВрд▓ рдХреБрдВрдЬреА рдЦрд╛рд▓реА рдереА рдпрд╛ рд╕рдореНрдорд┐рд▓рд┐рдд рдХреБрдВрдЬреА рд╕реЗ рдореЗрд▓ рдЦрд╛рддреА рдереА, рддреЛ рдХреЛрдб рдХреЛ рд╕рдореНрдорд┐рд▓рди рдХреЗ рд▓рд┐рдП рдПрдХ рдЙрдкрдпреБрдХреНрдд рд╕реНрд▓реЙрдЯ рдорд┐рд▓рд╛ рдФрд░ рд╕рдореНрдорд┐рд▓рд┐рдд рдореВрд▓реНрдп рдХреЛ рд╕реНрд▓реЙрдЯ рдореЗрдВ рдбрд╛рд▓рд╛ рдЧрдпрд╛ред

рдпрджрд┐ рдПрдХ рдХрд░реНрдиреЗрд▓ рдХреЙрд▓ рдореЗрдВ gpu_hashtable_insert() рдПрдХ рд╣реА рдХреБрдВрдЬреА рдХреЗ рд╕рд╛рде рдХрдИ рддрддреНрд╡ рд╣реИрдВ, рддреЛ рдЙрдирдХреЗ рдХрд┐рд╕реА рднреА рдорд╛рди рдХреЛ рдХреБрдВрдЬреА рд╕реНрд▓реЙрдЯ рдореЗрдВ рд▓рд┐рдЦрд╛ рдЬрд╛ рд╕рдХрддрд╛ рд╣реИред рдЗрд╕реЗ рд╕рд╛рдорд╛рдиреНрдп рдорд╛рдирд╛ рдЬрд╛рддрд╛ рд╣реИ: рдХреЙрд▓ рдХреЗ рджреМрд░рд╛рди рдХреА-рд╡реИрд▓реНрдпреВ рд░рд╛рдЗрдЯреНрд╕ рдореЗрдВ рд╕реЗ рдПрдХ рд╕рдлрд▓ рд╣реЛрдЧрд╛, рд▓реЗрдХрд┐рди рдЪреВрдВрдХрд┐ рдпрд╣ рд╕рдм рдирд┐рд╖реНрдкрд╛рджрди рдХреЗ рдХрдИ рдереНрд░реЗрдбреНрд╕ рдХреЗ рднреАрддрд░ рд╕рдорд╛рдирд╛рдВрддрд░ рдореЗрдВ рд╣реЛрддрд╛ рд╣реИ, рдЗрд╕рд▓рд┐рдП рд╣рдо рдпрд╣ рдЕрдиреБрдорд╛рди рдирд╣реАрдВ рд▓рдЧрд╛ рд╕рдХрддреЗ рдХрд┐ рдХреМрди рд╕реА рдореЗрдореЛрд░реА рд░рд╛рдЗрдЯрд┐рдВрдЧ рдЖрдЦрд┐рд░реА рд╣реЛрдЧреАред

рд╣реИрд╢ рдЯреЗрдмрд▓ рд▓реБрдХрдЕрдк

рдХреБрдВрдЬреА рдЦреЛрдЬрдиреЗ рдХреЗ рд▓рд┐рдП рдХреЛрдб:

uint32_t gpu_hashtable_lookup(KeyValue* hashtable, uint32_t key)
{
        uint32_t slot = hash(key);

        while (true)
        {
            if (hashtable[slot].key == key)
            {
                return hashtable[slot].value;
            }
            if (hashtable[slot].key == kEmpty)
            {
                return kEmpty;
            }
            slot = (slot + 1) & (kHashTableCapacity - 1);
        }
}

рдХрд┐рд╕реА рддрд╛рд▓рд┐рдХрд╛ рдореЗрдВ рд╕рдВрдЧреНрд░рд╣реАрдд рдХреБрдВрдЬреА рдХрд╛ рдорд╛рди рдЬреНрдЮрд╛рдд рдХрд░рдиреЗ рдХреЗ рд▓рд┐рдП, рд╣рдо рдЬрд┐рд╕ рдХреБрдВрдЬреА рдХреА рддрд▓рд╛рд╢ рдХрд░ рд░рд╣реЗ рд╣реИрдВ рдЙрд╕рдХреЗ рд╣реИрд╢ рд╕реЗ рд╢реБрд░реВ рдХрд░рдХреЗ рд╕рд░рдгреА рдХреЗ рдорд╛рдзреНрдпрдо рд╕реЗ рдкреБрдирд░рд╛рд╡реГрддрд┐ рдХрд░рддреЗ рд╣реИрдВред рдкреНрд░рддреНрдпреЗрдХ рд╕реНрд▓реЙрдЯ рдореЗрдВ, рд╣рдо рдЬрд╛рдВрдЪрддреЗ рд╣реИрдВ рдХрд┐ рдХреНрдпрд╛ рдХреБрдВрдЬреА рд╡рд╣ рд╣реИ рдЬрд┐рд╕реЗ рд╣рдо рдвреВрдВрдв рд░рд╣реЗ рд╣реИрдВ, рдФрд░ рдпрджрд┐ рд╣рд╛рдВ, рддреЛ рд╣рдо рдЙрд╕рдХрд╛ рдорд╛рди рд▓реМрдЯрд╛рддреЗ рд╣реИрдВред рд╣рдо рдпрд╣ рднреА рдЬрд╛рдВрдЪрддреЗ рд╣реИрдВ рдХрд┐ рдХреБрдВрдЬреА рдЦрд╛рд▓реА рд╣реИ рдпрд╛ рдирд╣реАрдВ, рдФрд░ рдпрджрд┐ рд╣рд╛рдВ, рддреЛ рд╣рдо рдЦреЛрдЬ рдХреЛ рд░реЛрдХ рджреЗрддреЗ рд╣реИрдВред

рдпрджрд┐ рд╣рдореЗрдВ рдХреБрдВрдЬреА рдирд╣реАрдВ рдорд┐рд▓ рдкрд╛рддреА рд╣реИ, рддреЛ рдХреЛрдб рдПрдХ рдЦрд╛рд▓реА рдорд╛рди рд▓реМрдЯрд╛рддрд╛ рд╣реИред

рдЗрди рд╕рднреА рдЦреЛрдЬ рдХрд╛рд░реНрдпреЛрдВ рдХреЛ рд╕рдореНрдорд┐рд▓рди рдФрд░ рд╡рд┐рд▓реЛрдкрди рдХреЗ рдорд╛рдзреНрдпрдо рд╕реЗ рдПрдХ рд╕рд╛рде рдирд┐рд╖реНрдкрд╛рджрд┐рдд рдХрд┐рдпрд╛ рдЬрд╛ рд╕рдХрддрд╛ рд╣реИред рддрд╛рд▓рд┐рдХрд╛ рдореЗрдВ рдкреНрд░рддреНрдпреЗрдХ рдЬреЛрдбрд╝реА рдореЗрдВ рдкреНрд░рд╡рд╛рд╣ рдХреЗ рд▓рд┐рдП рдКрдкрд░ рд╡рд░реНрдгрд┐рдд рдЪрд╛рд░ рд╕реНрдерд┐рддрд┐рдпреЛрдВ рдореЗрдВ рд╕реЗ рдПрдХ рд╣реЛрдЧреАред

рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдореЗрдВ рд╣рдЯрд╛рдирд╛

рдХреБрдВрдЬрд┐рдпрд╛рдБ рд╣рдЯрд╛рдиреЗ рдХреЗ рд▓рд┐рдП рдХреЛрдб:

void gpu_hashtable_delete(KeyValue* hashtable, uint32_t key, uint32_t value)
{
    uint32_t slot = hash(key);

    while (true)
    {
        if (hashtable[slot].key == key)
        {
            hashtable[slot].value = kEmpty;
            return;
        }
        if (hashtable[slot].key == kEmpty)
        {
            return;
        }
        slot = (slot + 1) & (kHashTableCapacity - 1);
    }
}

рдХрд┐рд╕реА рдХреБрдВрдЬреА рдХреЛ рд╣рдЯрд╛рдирд╛ рдПрдХ рдЕрд╕рд╛рдорд╛рдиреНрдп рддрд░реАрдХреЗ рд╕реЗ рдХрд┐рдпрд╛ рдЬрд╛рддрд╛ рд╣реИ: рд╣рдо рдХреБрдВрдЬреА рдХреЛ рддрд╛рд▓рд┐рдХрд╛ рдореЗрдВ рдЫреЛрдбрд╝ рджреЗрддреЗ рд╣реИрдВ рдФрд░ рдЙрд╕рдХреЗ рдорд╛рди (рдХреБрдВрдЬреА рдХреЛ рдирд╣реАрдВ) рдХреЛ рдЦрд╛рд▓реА рдХреЗ рд░реВрдк рдореЗрдВ рдЪрд┐рд╣реНрдирд┐рдд рдХрд░рддреЗ рд╣реИрдВред рдпрд╣ рдХреЛрдб рдХрд╛рдлреА рд╣рдж рддрдХ рдорд┐рд▓рддрд╛ рдЬреБрд▓рддрд╛ рд╣реИ lookup(), рд╕рд┐рд╡рд╛рдп рдЗрд╕рдХреЗ рдХрд┐ рдЬрдм рдХрд┐рд╕реА рдХреБрдВрдЬреА рдкрд░ рдорд┐рд▓рд╛рди рдкрд╛рдпрд╛ рдЬрд╛рддрд╛ рд╣реИ, рддреЛ рдпрд╣ рдЙрд╕рдХрд╛ рдорд╛рди рд░рд┐рдХреНрдд рдХрд░ рджреЗрддрд╛ рд╣реИред

рдЬреИрд╕рд╛ рдХрд┐ рдКрдкрд░ рдЙрд▓реНрд▓реЗрдЦ рдХрд┐рдпрд╛ рдЧрдпрд╛ рд╣реИ, рдПрдХ рдмрд╛рд░ рдПрдХ рд╕реНрд▓реЙрдЯ рдореЗрдВ рдХреБрдВрдЬреА рд▓рд┐рдЦреА рдЬрд╛рдиреЗ рдХреЗ рдмрд╛рдж, рдЗрд╕реЗ рд╕реНрдерд╛рдирд╛рдВрддрд░рд┐рдд рдирд╣реАрдВ рдХрд┐рдпрд╛ рдЬрд╛рддрд╛ рд╣реИред рдпрд╣рд╛рдВ рддрдХ тАЛтАЛрдХрд┐ рдЬрдм рдХреЛрдИ рддрддреНрд╡ рддрд╛рд▓рд┐рдХрд╛ рд╕реЗ рд╣рдЯрд╛ рджрд┐рдпрд╛ рдЬрд╛рддрд╛ рд╣реИ, рддрдм рднреА рдХреБрдВрдЬреА рдпрдерд╛рд╡рдд рд░рд╣рддреА рд╣реИ, рдЙрд╕рдХрд╛ рдорд╛рди рдмрд╕ рдЦрд╛рд▓реА рд╣реЛ рдЬрд╛рддрд╛ рд╣реИред рдЗрд╕рдХрд╛ рдорддрд▓рдм рдпрд╣ рд╣реИ рдХрд┐ рд╣рдореЗрдВ рд╕реНрд▓реЙрдЯ рдорд╛рди рдХреЗ рд▓рд┐рдП рдкрд░рдорд╛рдгреБ рд▓реЗрдЦрди рдСрдкрд░реЗрд╢рди рдХрд╛ рдЙрдкрдпреЛрдЧ рдХрд░рдиреЗ рдХреА рдЖрд╡рд╢реНрдпрдХрддрд╛ рдирд╣реАрдВ рд╣реИ, рдХреНрдпреЛрдВрдХрд┐ рдЗрд╕рд╕реЗ рдХреЛрдИ рдлрд░реНрдХ рдирд╣реАрдВ рдкрдбрд╝рддрд╛ рдХрд┐ рд╡рд░реНрддрдорд╛рди рдорд╛рди рдЦрд╛рд▓реА рд╣реИ рдпрд╛ рдирд╣реАрдВ - рдпрд╣ рдЕрднреА рднреА рдЦрд╛рд▓реА рд╣реЛ рдЬрд╛рдПрдЧрд╛ред

рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдХрд╛ рдЖрдХрд╛рд░ рдмрджрд▓рдирд╛

рдЖрдк рдПрдХ рдмрдбрд╝реА рддрд╛рд▓рд┐рдХрд╛ рдмрдирд╛рдХрд░ рдФрд░ рдЙрд╕рдореЗрдВ рдкреБрд░рд╛рдиреА рддрд╛рд▓рд┐рдХрд╛ рд╕реЗ рдЧреИрд░-рд░рд┐рдХреНрдд рддрддреНрд╡реЛрдВ рдХреЛ рд╕рдореНрдорд┐рд▓рд┐рдд рдХрд░рдХреЗ рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдХрд╛ рдЖрдХрд╛рд░ рдмрджрд▓ рд╕рдХрддреЗ рд╣реИрдВред рдореИрдВрдиреЗ рдЗрд╕ рдХрд╛рд░реНрдпрдХреНрд╖рдорддрд╛ рдХреЛ рд▓рд╛рдЧреВ рдирд╣реАрдВ рдХрд┐рдпрд╛ рдХреНрдпреЛрдВрдХрд┐ рдореИрдВ рдирдореВрдирд╛ рдХреЛрдб рдХреЛ рд╕рд░рд▓ рд░рдЦрдирд╛ рдЪрд╛рд╣рддрд╛ рдерд╛ред рдЗрд╕рдХреЗ рдЕрд▓рд╛рд╡рд╛, CUDA рдкреНрд░реЛрдЧреНрд░рд╛рдореЛрдВ рдореЗрдВ, рдореЗрдореЛрд░реА рдЖрд╡рдВрдЯрди рдЕрдХреНрд╕рд░ CUDA рдХрд░реНрдиреЗрд▓ рдХреЗ рдмрдЬрд╛рдп рд╣реЛрд╕реНрдЯ рдХреЛрдб рдореЗрдВ рдХрд┐рдпрд╛ рдЬрд╛рддрд╛ рд╣реИред

рд▓реЗрдЦ рдПрдХ рд▓реЙрдХ-рдореБрдХреНрдд рдкреНрд░рддреАрдХреНрд╖рд╛-рдореБрдХреНрдд рд╣реИрд╢ рдЯреЗрдмрд▓ рд╡рд░реНрдгрди рдХрд░рддрд╛ рд╣реИ рдХрд┐ рдРрд╕реА рд▓реЙрдХ-рд╕рдВрд░рдХреНрд╖рд┐рдд рдбреЗрдЯрд╛ рд╕рдВрд░рдЪрдирд╛ рдХреЛ рдХреИрд╕реЗ рд╕рдВрд╢реЛрдзрд┐рдд рдХрд┐рдпрд╛ рдЬрд╛рдПред

рдкреНрд░рддрд┐рд╕реНрдкрд░реНрдзрд╛

рдЙрдкрд░реЛрдХреНрдд рдлрд╝рдВрдХреНрд╢рди рдХреЛрдб рд╕реНрдирд┐рдкреЗрдЯреНрд╕ рдореЗрдВ gpu_hashtable_insert(), _lookup() ╨╕ _delete() рдПрдХ рд╕рдордп рдореЗрдВ рдПрдХ рдХреБрдВрдЬреА-рдореВрд▓реНрдп рдЬреЛрдбрд╝реА рдХреЛ рд╕рдВрд╕рд╛рдзрд┐рдд рдХрд░реЗрдВред рдФрд░ рдирд┐рдЪрд▓рд╛ gpu_hashtable_insert(), _lookup() ╨╕ _delete() рдЬреЛрдбрд╝реЛрдВ рдХреА рдПрдХ рд╕рд░рдгреА рдХреЛ рд╕рдорд╛рдирд╛рдВрддрд░ рдореЗрдВ рд╕рдВрд╕рд╛рдзрд┐рдд рдХрд░реЗрдВ, рдкреНрд░рддреНрдпреЗрдХ рдЬреЛрдбрд╝реА рдХреЛ рдПрдХ рдЕрд▓рдЧ GPU рдирд┐рд╖реНрдкрд╛рджрди рдереНрд░реЗрдб рдореЗрдВ:

// CPU code to invoke the CUDA kernel on the GPU
uint32_t threadblocksize = 1024;
uint32_t gridsize = (numkvs + threadblocksize - 1) / threadblocksize;
gpu_hashtable_insert_kernel<<<gridsize, threadblocksize>>>(hashtable, kvs, numkvs);

// GPU code to process numkvs key/values in parallel
void gpu_hashtable_insert_kernel(KeyValue* hashtable, const KeyValue* kvs, unsigned int numkvs)
{
    unsigned int threadid = blockIdx.x*blockDim.x + threadIdx.x;
    if (threadid < numkvs)
    {
        gpu_hashtable_insert(hashtable, kvs[threadid].key, kvs[threadid].value);
    }
}

рд▓реЙрдХ-рдкреНрд░рддрд┐рд░реЛрдзреА рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рд╕рдорд╡рд░реНрддреА рдкреНрд░рд╡рд┐рд╖реНрдЯрд┐рдпреЛрдВ, рд▓реБрдХрдЕрдк рдФрд░ рд╡рд┐рд▓реЛрдкрди рдХрд╛ рд╕рдорд░реНрдерди рдХрд░рддреА рд╣реИред рдЪреВрдБрдХрд┐ рдХреБрдВрдЬреА-рдореВрд▓реНрдп рдЬреЛрдбрд╝реЗ рд╣рдореЗрд╢рд╛ рдЪрд╛рд░ рдЕрд╡рд╕реНрдерд╛рдУрдВ рдореЗрдВ рд╕реЗ рдПрдХ рдореЗрдВ рд╣реЛрддреЗ рд╣реИрдВ рдФрд░ рдХреБрдВрдЬрд┐рдпрд╛рдБ рд╣рд┐рд▓рддреА рдирд╣реАрдВ рд╣реИрдВ, рддрд╛рд▓рд┐рдХрд╛ рддрдм рднреА рд╢реБрджреНрдзрддрд╛ рдХреА рдЧрд╛рд░рдВрдЯреА рджреЗрддреА рд╣реИ рдЬрдм рд╡рд┐рднрд┐рдиреНрди рдкреНрд░рдХрд╛рд░ рдХреЗ рд╕рдВрдЪрд╛рд▓рди рдПрдХ рд╕рд╛рде рдЙрдкрдпреЛрдЧ рдХрд┐рдП рдЬрд╛рддреЗ рд╣реИрдВред

рд╣рд╛рд▓рд╛рдБрдХрд┐, рдпрджрд┐ рд╣рдо рд╕рдореНрдорд┐рд▓рди рдФрд░ рд╡рд┐рд▓реЛрдкрди рдХреЗ рдПрдХ рдмреИрдЪ рдХреЛ рд╕рдорд╛рдирд╛рдВрддрд░ рдореЗрдВ рд╕рдВрд╕рд╛рдзрд┐рдд рдХрд░рддреЗ рд╣реИрдВ, рдФрд░ рдпрджрд┐ рдЬреЛрдбрд╝реЗ рдХреЗ рдЗрдирдкреБрдЯ рд╕рд░рдгреА рдореЗрдВ рдбреБрдкреНрд▓рд┐рдХреЗрдЯ рдХреБрдВрдЬрд┐рдпрд╛рдБ рд╣реИрдВ, рддреЛ рд╣рдо рдпрд╣ рдЕрдиреБрдорд╛рди рдирд╣реАрдВ рд▓рдЧрд╛ рдкрд╛рдПрдВрдЧреЗ рдХрд┐ рдХреМрди рд╕реА рдЬреЛрдбрд╝реА "рдЬреАрддреЗрдЧреА" - рдЬреЛ рдХрд┐ рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдореЗрдВ рд╕рдмрд╕реЗ рдЕрдВрдд рдореЗрдВ рд▓рд┐рдЦреА рдЬрд╛рдПрдЧреАред рдорд╛рди рд▓реАрдЬрд┐рдП рдХрд┐ рд╣рдордиреЗ рдЬреЛрдбрд╝реЗ рдХреЗ рдЗрдирдкреБрдЯ рдРрд░реЗ рдХреЗ рд╕рд╛рде рдЗрдВрд╕рд░реНрд╢рди рдХреЛрдб рдХреЛ рдХреЙрд▓ рдХрд┐рдпрд╛ рд╣реИ A/0 B/1 A/2 C/3 A/4. рдЬрдм рдХреЛрдб рдкреВрд░рд╛ рд╣реЛ рдЬрд╛рдП, рдЬреЛрдбрд╝реЗ B/1 ╨╕ C/3 рддрд╛рд▓рд┐рдХрд╛ рдореЗрдВ рдореМрдЬреВрдж рд╣реЛрдиреЗ рдХреА рдЧрд╛рд░рдВрдЯреА рд╣реИ, рд▓реЗрдХрд┐рди рд╕рд╛рде рд╣реА рдХреЛрдИ рднреА рдЬреЛрдбрд╝рд╛ рдЗрд╕рдореЗрдВ рджрд┐рдЦрд╛рдИ рджреЗрдЧрд╛ A/0, A/2 рдпрд╛ A/4. рдпрд╣ рдХреЛрдИ рд╕рдорд╕реНрдпрд╛ рд╣реЛ рднреА рд╕рдХрддреА рд╣реИ рдФрд░ рдирд╣реАрдВ рднреА - рдпрд╣ рд╕рдм рдПрдкреНрд▓рд┐рдХреЗрд╢рди рдкрд░ рдирд┐рд░реНрднрд░ рдХрд░рддрд╛ рд╣реИред рдЖрдк рдкрд╣рд▓реЗ рд╕реЗ рдЬрд╛рди рд╕рдХрддреЗ рд╣реИрдВ рдХрд┐ рдЗрдирдкреБрдЯ рдРрд░реЗ рдореЗрдВ рдХреЛрдИ рдбреБрдкреНрд▓рд┐рдХреЗрдЯ рдХреБрдВрдЬрд┐рдпрд╛рдБ рдирд╣реАрдВ рд╣реИрдВ, рдпрд╛ рдЖрдкрдХреЛ рдЗрд╕ рдмрд╛рдд рдХреА рдкрд░рд╡рд╛рд╣ рдирд╣реАрдВ рд╣реЛрдЧреА рдХрд┐ рдЕрдВрддрд┐рдо рдмрд╛рд░ рдХреМрди рд╕рд╛ рдорд╛рди рд▓рд┐рдЦрд╛ рдЧрдпрд╛ рдерд╛ред

рдпрджрд┐ рдпрд╣ рдЖрдкрдХреЗ рд▓рд┐рдП рдПрдХ рд╕рдорд╕реНрдпрд╛ рд╣реИ, рддреЛ рдЖрдкрдХреЛ рдбреБрдкреНрд▓рд┐рдХреЗрдЯ рдЬреЛрдбрд╝рд┐рдпреЛрдВ рдХреЛ рдЕрд▓рдЧ-рдЕрд▓рдЧ CUDA рд╕рд┐рд╕реНрдЯрдо рдХреЙрд▓ рдореЗрдВ рдЕрд▓рдЧ рдХрд░рдирд╛ рд╣реЛрдЧрд╛ред CUDA рдореЗрдВ, рдХрд░реНрдиреЗрд▓ рдХреЛ рдХреЙрд▓ рдХрд░рдиреЗ рд╡рд╛рд▓рд╛ рдХреЛрдИ рднреА рдСрдкрд░реЗрд╢рди рд╣рдореЗрд╢рд╛ рдЕрдЧрд▓реЗ рдХрд░реНрдиреЗрд▓ рдХреЙрд▓ рд╕реЗ рдкрд╣рд▓реЗ рдкреВрд░рд╛ рд╣реЛрддрд╛ рд╣реИ (рдХрдо рд╕реЗ рдХрдо рдПрдХ рдереНрд░реЗрдб рдХреЗ рднреАрддрд░ред рд╡рд┐рднрд┐рдиреНрди рдереНрд░реЗрдб рдореЗрдВ, рдХрд░реНрдиреЗрд▓ рдХреЛ рд╕рдорд╛рдирд╛рдВрддрд░ рдореЗрдВ рдирд┐рд╖реНрдкрд╛рджрд┐рдд рдХрд┐рдпрд╛ рдЬрд╛рддрд╛ рд╣реИ)ред рдЙрдкрд░реЛрдХреНрдд рдЙрджрд╛рд╣рд░рдг рдореЗрдВ, рдпрджрд┐ рдЖрдк рдПрдХ рдХрд░реНрдиреЗрд▓ рдХреЛ рдХреЙрд▓ рдХрд░рддреЗ рд╣реИрдВ A/0 B/1 A/2 C/3, рдФрд░ рджреВрд╕рд░реЗ рдХреЗ рд╕рд╛рде A/4, рдлрд┐рд░ рдХреБрдВрдЬреА A рдореВрд▓реНрдп рдорд┐рд▓реЗрдЧрд╛ 4.

рдЕрдм рдмрд╛рдд рдХрд░рддреЗ рд╣реИрдВ рдХрд┐ рдХреНрдпрд╛ рдХрд╛рд░реНрдпреЛрдВ рдХреЛ рдХрд░рдирд╛ рдЪрд╛рд╣рд┐рдП lookup() ╨╕ delete() рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдореЗрдВ рдЬреЛрдбрд╝реЗ рдХреА рдПрдХ рд╕рд░рдгреА рдХреЗ рд▓рд┐рдП рдПрдХ рд╕рд╛рджреЗ рдпрд╛ рдЕрд╕реНрдерд┐рд░ рд╕реВрдЪрдХ рдХрд╛ рдЙрдкрдпреЛрдЧ рдХрд░реЗрдВред рд╕реАрдпреВрдбреАрдП рджрд╕реНрддрд╛рд╡реЗрдЬрд╝реАрдХрд░рдг рдмрддрд╛рддрд╛ рд╣реИ:

рдХрдВрдкрд╛рдЗрд▓рд░ рдкрдврд╝рдиреЗ рдФрд░ рд▓рд┐рдЦрдиреЗ рдХреЛ рд╡реИрд╢реНрд╡рд┐рдХ рдпрд╛ рд╕рд╛рдЭрд╛ рдореЗрдореЛрд░реА рдореЗрдВ рдЕрдиреБрдХреВрд▓рд┐рдд рдХрд░рдирд╛ рдЪреБрди рд╕рдХрддрд╛ рд╣реИ... рдЗрди рдЕрдиреБрдХреВрд▓рди рдХреЛ рдХреАрд╡рд░реНрдб рдХрд╛ рдЙрдкрдпреЛрдЧ рдХрд░рдХреЗ рдЕрдХреНрд╖рдо рдХрд┐рдпрд╛ рдЬрд╛ рд╕рдХрддрд╛ рд╣реИ volatile: ...рдЗрд╕ рдЪрд░ рдХрд╛ рдХреЛрдИ рднреА рд╕рдВрджрд░реНрдн рд╡рд╛рд╕реНрддрд╡рд┐рдХ рдореЗрдореЛрд░реА рдореЗрдВ рдкрдврд╝рдиреЗ рдпрд╛ рд▓рд┐рдЦрдиреЗ рдХреЗ рдирд┐рд░реНрджреЗрд╢ рдореЗрдВ рд╕рдВрдХрд▓рд┐рдд рдХрд┐рдпрд╛ рдЬрд╛рддрд╛ рд╣реИред

рд╢реБрджреНрдзрддрд╛ рд╕рдВрдмрдВрдзреА рд╡рд┐рдЪрд╛рд░реЛрдВ рдХреЗ рд▓рд┐рдП рдЖрд╡реЗрджрди рдХреА рдЖрд╡рд╢реНрдпрдХрддрд╛ рдирд╣реАрдВ рд╣реЛрддреА рд╣реИ volatile. рдпрджрд┐ рдирд┐рд╖реНрдкрд╛рджрди рдереНрд░реЗрдб рдкрд╣рд▓реЗ рдкрдврд╝реЗ рдЧрдП рдСрдкрд░реЗрд╢рди рд╕реЗ рдХреИрд╢реНрдб рдорд╛рди рдХрд╛ рдЙрдкрдпреЛрдЧ рдХрд░рддрд╛ рд╣реИ, рддреЛ рдпрд╣ рдереЛрдбрд╝реА рдкреБрд░рд╛рдиреА рдЬрд╛рдирдХрд╛рд░реА рдХрд╛ рдЙрдкрдпреЛрдЧ рдХрд░реЗрдЧрд╛ред рд▓реЗрдХрд┐рди рдлрд┐рд░ рднреА, рдпрд╣ рдХрд░реНрдиреЗрд▓ рдХреЙрд▓ рдХреЗ рдПрдХ рдирд┐рд╢реНрдЪрд┐рдд рдХреНрд╖рдг рдореЗрдВ рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдХреА рд╕рд╣реА рд╕реНрдерд┐рддрд┐ рдХреА рдЬрд╛рдирдХрд╛рд░реА рд╣реИред рдпрджрд┐ рдЖрдкрдХреЛ рдирд╡реАрдирддрдо рдЬрд╛рдирдХрд╛рд░реА рдХрд╛ рдЙрдкрдпреЛрдЧ рдХрд░рдиреЗ рдХреА рдЖрд╡рд╢реНрдпрдХрддрд╛ рд╣реИ, рддреЛ рдЖрдк рд╕реВрдЪрдХрд╛рдВрдХ рдХрд╛ рдЙрдкрдпреЛрдЧ рдХрд░ рд╕рдХрддреЗ рд╣реИрдВ volatile, рд▓реЗрдХрд┐рди рддрдм рдкреНрд░рджрд░реНрд╢рди рдереЛрдбрд╝рд╛ рдХрдо рд╣реЛ рдЬрд╛рдПрдЧрд╛: рдореЗрд░реЗ рдкрд░реАрдХреНрд╖рдгреЛрдВ рдХреЗ рдЕрдиреБрд╕рд╛рд░, 32 рдорд┐рд▓рд┐рдпрди рддрддреНрд╡реЛрдВ рдХреЛ рд╣рдЯрд╛рдиреЗ рдкрд░, рдЧрддрд┐ 500 тАЛтАЛрдорд┐рд▓рд┐рдпрди рд╡рд┐рд▓реЛрдкрди/рд╕реЗрдХрдВрдб рд╕реЗ рдШрдЯрдХрд░ 450 рдорд┐рд▓рд┐рдпрди рд╡рд┐рд▓реЛрдкрди/рд╕реЗрдХрдВрдб рд╣реЛ рдЧрдИред

рдирд┐рд╖реНрдкрд╛рджрди

64 рдорд┐рд▓рд┐рдпрди рддрддреНрд╡реЛрдВ рдХреЛ рд╕рдореНрдорд┐рд▓рд┐рдд рдХрд░рдиреЗ рдФрд░ рдЙрдирдореЗрдВ рд╕реЗ 32 рдорд┐рд▓рд┐рдпрди рдХреЛ рд╣рдЯрд╛рдиреЗ рдХреЗ рдкрд░реАрдХреНрд╖рдг рдореЗрдВ, рдХреЗ рдмреАрдЪ рдкреНрд░рддрд┐рд╕реНрдкрд░реНрдзрд╛ std::unordered_map рдФрд░ GPU рдХреЗ рд▓рд┐рдП рд╡рд╕реНрддреБрддрдГ рдХреЛрдИ рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдирд╣реАрдВ рд╣реИ:

GPU рдХреЗ рд▓рд┐рдП рд╕рд░рд▓ рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛
std::unordered_map рддрддреНрд╡реЛрдВ рдХреЛ рдбрд╛рд▓рдиреЗ рдФрд░ рд╣рдЯрд╛рдиреЗ рдФрд░ рдлрд┐рд░ рдЙрдиреНрд╣реЗрдВ рдореБрдХреНрдд рдХрд░рдиреЗ рдореЗрдВ 70 рдПрдордПрд╕ рдЦрд░реНрдЪ рд╣реБрдП unordered_map (рд▓рд╛рдЦреЛрдВ рддрддреНрд╡реЛрдВ рд╕реЗ рдЫреБрдЯрдХрд╛рд░рд╛ рдкрд╛рдиреЗ рдореЗрдВ рдмрд╣реБрдд рд╕рдордп рд▓рдЧрддрд╛ рд╣реИ, рдХреНрдпреЛрдВрдХрд┐ рдЕрдВрджрд░ unordered_map рдПрдХрд╛рдзрд┐рдХ рдореЗрдореЛрд░реА рдЖрд╡рдВрдЯрди рдХрд┐рдП рдЬрд╛рддреЗ рд╣реИрдВ)ред рд╕рдЪ рдмреЛрд▓реВ рддреЛ, std:unordered_map рдкреВрд░реА рддрд░рд╣ рд╕реЗ рдЕрд▓рдЧ рдкреНрд░рддрд┐рдмрдВрдз. рдпрд╣ рдирд┐рд╖реНрдкрд╛рджрди рдХрд╛ рдПрдХ рдПрдХрд▓ рд╕реАрдкреАрдпреВ рдереНрд░реЗрдб рд╣реИ, рдХрд┐рд╕реА рднреА рдЖрдХрд╛рд░ рдХреЗ рдХреБрдВрдЬреА-рдорд╛рдиреЛрдВ рдХрд╛ рд╕рдорд░реНрдерди рдХрд░рддрд╛ рд╣реИ, рдЙрдЪреНрдЪ рдЙрдкрдпреЛрдЧ рджрд░ рдкрд░ рдЕрдЪреНрдЫрд╛ рдкреНрд░рджрд░реНрд╢рди рдХрд░рддрд╛ рд╣реИ, рдФрд░ рдХрдИ рд╡рд┐рд▓реЛрдкрди рдХреЗ рдмрд╛рдж рд╕реНрдерд┐рд░ рдкреНрд░рджрд░реНрд╢рди рджрд┐рдЦрд╛рддрд╛ рд╣реИред

GPU рдФрд░ рдЕрдВрддрд░-рдкреНрд░реЛрдЧреНрд░рд╛рдо рд╕рдВрдЪрд╛рд░ рдХреЗ рд▓рд┐рдП рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдХреА рдЕрд╡рдзрд┐ 984 рдПрдордПрд╕ рдереАред рдЗрд╕рдореЗрдВ рддрд╛рд▓рд┐рдХрд╛ рдХреЛ рдореЗрдореЛрд░реА рдореЗрдВ рд░рдЦрдиреЗ рдФрд░ рдЙрд╕реЗ рд╣рдЯрд╛рдиреЗ (рдПрдХ рдмрд╛рд░ рдореЗрдВ 1 рдЬреАрдмреА рдореЗрдореЛрд░реА рдЖрд╡рдВрдЯрд┐рдд рдХрд░рдирд╛, рдЬрд┐рд╕рдореЗрдВ CUDA рдореЗрдВ рдХреБрдЫ рд╕рдордп рд▓рдЧрддрд╛ рд╣реИ), рддрддреНрд╡реЛрдВ рдХреЛ рд╕рдореНрдорд┐рд▓рд┐рдд рдХрд░рдирд╛ рдФрд░ рд╣рдЯрд╛рдирд╛, рдФрд░ рдЙрди рдкрд░ рдкреБрдирд░рд╛рд╡реГрддреНрддрд┐ рдХрд░рдиреЗ рдореЗрдВ рд▓рдЧрдиреЗ рд╡рд╛рд▓рд╛ рд╕рдордп рд╢рд╛рдорд┐рд▓ рд╣реИред рд╡реАрдбрд┐рдпреЛ рдХрд╛рд░реНрдб рдореЗрдореЛрд░реА рд╕реЗ рдЖрдиреЗ рдФрд░ рдЬрд╛рдиреЗ рд╡рд╛рд▓реА рд╕рднреА рдкреНрд░рддрд┐рдпреЛрдВ рдХреЛ рднреА рдзреНрдпрд╛рди рдореЗрдВ рд░рдЦрд╛ рдЬрд╛рддрд╛ рд╣реИред

рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдХреЛ рдкреВрд░рд╛ рд╣реЛрдиреЗ рдореЗрдВ 271 рдПрдордПрд╕ рд▓рдЧреЗред рдЗрд╕рдореЗрдВ рд╡реАрдбрд┐рдпреЛ рдХрд╛рд░реНрдб рджреНрд╡рд╛рд░рд╛ рддрддреНрд╡реЛрдВ рдХреЛ рдбрд╛рд▓рдиреЗ рдФрд░ рд╣рдЯрд╛рдиреЗ рдореЗрдВ рдмрд┐рддрд╛рдпрд╛ рдЧрдпрд╛ рд╕рдордп рд╢рд╛рдорд┐рд▓ рд╣реИ, рдФрд░ рдореЗрдореЛрд░реА рдореЗрдВ рдХреЙрдкреА рдХрд░рдиреЗ рдФрд░ рдкрд░рд┐рдгрд╛рдореА рддрд╛рд▓рд┐рдХрд╛ рдкрд░ рдкреБрдирд░рд╛рд╡реГрддреНрдд рдХрд░рдиреЗ рдореЗрдВ рд▓рдЧрдиреЗ рд╡рд╛рд▓реЗ рд╕рдордп рдХреЛ рдзреНрдпрд╛рди рдореЗрдВ рдирд╣реАрдВ рд░рдЦрд╛ рдЬрд╛рддрд╛ рд╣реИред рдпрджрд┐ GPU рддрд╛рд▓рд┐рдХрд╛ рд▓рдВрдмреЗ рд╕рдордп рддрдХ рд░рд╣рддреА рд╣реИ, рдпрд╛ рдпрджрд┐ рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдкреВрд░реА рддрд░рд╣ рд╕реЗ рд╡реАрдбрд┐рдпреЛ рдХрд╛рд░реНрдб рдХреА рдореЗрдореЛрд░реА рдореЗрдВ рд╕рдорд╛рд╣рд┐рдд рд╣реИ (рдЙрджрд╛рд╣рд░рдг рдХреЗ рд▓рд┐рдП, рдПрдХ рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдмрдирд╛рдиреЗ рдХреЗ рд▓рд┐рдП рдЬрд┐рд╕рдХрд╛ рдЙрдкрдпреЛрдЧ рдЕрдиреНрдп GPU рдХреЛрдб рджреНрд╡рд╛рд░рд╛ рдХрд┐рдпрд╛ рдЬрд╛рдПрдЧрд╛, рди рдХрд┐ рдХреЗрдВрджреНрд░реАрдп рдкреНрд░реЛрд╕реЗрд╕рд░ рджреНрд╡рд╛рд░рд╛), рддреЛ рдкрд░реАрдХреНрд╖рд╛ рдкрд░рд┐рдгрд╛рдо рдкреНрд░рд╛рд╕рдВрдЧрд┐рдХ рд╣реИ.

рд╡реАрдбрд┐рдпреЛ рдХрд╛рд░реНрдб рдХреЗ рд▓рд┐рдП рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдЙрдЪреНрдЪ рдереНрд░реВрдкреБрдЯ рдФрд░ рд╕рдХреНрд░рд┐рдп рд╕рдорд╛рдирд╛рдВрддрд░реАрдХрд░рдг рдХреЗ рдХрд╛рд░рдг рдЙрдЪреНрдЪ рдкреНрд░рджрд░реНрд╢рди рдкреНрд░рджрд░реНрд╢рд┐рдд рдХрд░рддреА рд╣реИред

рд╕реАрдорд╛рдПрдВ

рд╣реИрд╢ рдЯреЗрдмрд▓ рдЖрд░реНрдХрд┐рдЯреЗрдХреНрдЪрд░ рдореЗрдВ рдХреБрдЫ рдореБрджреНрджреЛрдВ рдХреЗ рдмрд╛рд░реЗ рдореЗрдВ рдкрддрд╛ рд╣реЛрдирд╛ рдЪрд╛рд╣рд┐рдП:

  • рдХреНрд▓рд╕реНрдЯрд░рд┐рдВрдЧ рдХреЗ рдХрд╛рд░рдг рд░реИрдЦрд┐рдХ рдЬрд╛рдВрдЪ рдореЗрдВ рдмрд╛рдзрд╛ рдЖрддреА рд╣реИ, рдЬрд┐рд╕рдХреЗ рдХрд╛рд░рдг рддрд╛рд▓рд┐рдХрд╛ рдореЗрдВ рдХреБрдВрдЬрд┐рдпрд╛рдБ рдкреВрд░реА рддрд░рд╣ рд╕реЗ рдХрдо рд░рдЦреА рдЬрд╛рддреА рд╣реИрдВред
  • рдлрд╝рдВрдХреНрд╢рди рдХрд╛ рдЙрдкрдпреЛрдЧ рдХрд░рдХреЗ рдХреБрдВрдЬрд┐рдпрд╛рдБ рдирд╣реАрдВ рд╣рдЯрд╛рдИ рдЬрд╛рддреА рд╣реИрдВ delete рдФрд░ рд╕рдордп рдХреЗ рд╕рд╛рде рд╡реЗ рдореЗрдЬрд╝ рдХреЛ рдЕрд╡реНрдпрд╡рд╕реНрдерд┐рдд рдХрд░ рджреЗрддреЗ рд╣реИрдВред

рдкрд░рд┐рдгрд╛рдорд╕реНрд╡рд░реВрдк, рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдХрд╛ рдкреНрд░рджрд░реНрд╢рди рдзреАрд░реЗ-рдзреАрд░реЗ рдЦрд╝рд░рд╛рдм рд╣реЛ рд╕рдХрддрд╛ рд╣реИ, рдЦрд╛рд╕рдХрд░ рдпрджрд┐ рдпрд╣ рд▓рдВрдмреЗ рд╕рдордп рд╕реЗ рдореМрдЬреВрдж рд╣реИ рдФрд░ рдЗрд╕рдореЗрдВ рдХрдИ рдкреНрд░рд╡рд┐рд╖реНрдЯрд┐рдпрд╛рдБ рдФрд░ рд╡рд┐рд▓реЛрдкрди рд╣реИрдВред рдЗрди рдиреБрдХрд╕рд╛рдиреЛрдВ рдХреЛ рдХрдо рдХрд░рдиреЗ рдХрд╛ рдПрдХ рддрд░реАрдХрд╛ рдпрд╣ рд╣реИ рдХрд┐ рдХрд╛рдлреА рдХрдо рдЙрдкрдпреЛрдЧ рджрд░ рдХреЗ рд╕рд╛рде рдПрдХ рдирдИ рддрд╛рд▓рд┐рдХрд╛ рдореЗрдВ рджреЛрдмрд╛рд░рд╛ рдмрджрд▓рд╛рд╡ рдХрд┐рдпрд╛ рдЬрд╛рдП рдФрд░ рджреЛрдмрд╛рд░рд╛ рдЗрд╕реНрддреЗрдорд╛рд▓ рдХреЗ рджреМрд░рд╛рди рд╣рдЯрд╛рдИ рдЧрдИ рдХреБрдВрдЬрд┐рдпреЛрдВ рдХреЛ рдлрд╝рд┐рд▓реНрдЯрд░ рдХрд░ рджрд┐рдпрд╛ рдЬрд╛рдПред

рд╡рд░реНрдгрд┐рдд рдореБрджреНрджреЛрдВ рдХреЛ рд╕реНрдкрд╖реНрдЯ рдХрд░рдиреЗ рдХреЗ рд▓рд┐рдП, рдореИрдВ рдЙрдкрд░реЛрдХреНрдд рдХреЛрдб рдХрд╛ рдЙрдкрдпреЛрдЧ 128 рдорд┐рд▓рд┐рдпрди рддрддреНрд╡реЛрдВ рдХреЗ рд╕рд╛рде рдПрдХ рддрд╛рд▓рд┐рдХрд╛ рдмрдирд╛рдиреЗ рдФрд░ 4 рдорд┐рд▓рд┐рдпрди рддрддреНрд╡реЛрдВ рдХреЗ рдорд╛рдзреНрдпрдо рд╕реЗ рд▓реВрдк рдХрд░рдиреЗ рдХреЗ рд▓рд┐рдП рдХрд░реВрдВрдЧрд╛ рдЬрдм рддрдХ рдХрд┐ рдореИрдВ 124 рдорд┐рд▓рд┐рдпрди рд╕реНрд▓реЙрдЯ (рд▓рдЧрднрдЧ 0,96 рдХреА рдЙрдкрдпреЛрдЧ рджрд░) рдирд╣реАрдВ рднрд░ рд▓реЗрддрд╛ред рдпрд╣рд╛рдВ рдкрд░рд┐рдгрд╛рдо рддрд╛рд▓рд┐рдХрд╛ рд╣реИ, рдкреНрд░рддреНрдпреЗрдХ рдкрдВрдХреНрддрд┐ рдПрдХ рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдореЗрдВ 4 рдорд┐рд▓рд┐рдпрди рдирдП рддрддреНрд╡ рдбрд╛рд▓рдиреЗ рдХреЗ рд▓рд┐рдП рдПрдХ CUDA рдХрд░реНрдиреЗрд▓ рдХреЙрд▓ рд╣реИ:

рдЙрдкрдпреЛрдЧ рджрд░
рд╕рдореНрдорд┐рд▓рди рдЕрд╡рдзрд┐ 4 рддрддреНрд╡

0,00
11,608448 рдПрдордПрд╕ (361,314798 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,03
11,751424 рдПрдордПрд╕ (356,918799 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,06
11,942592 рдПрдордПрд╕ (351,205515 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,09
12,081120 рдПрдордПрд╕ (347,178429 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,12
12,242560 рдПрдордПрд╕ (342,600233 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,16
12,396448 рдПрдордПрд╕ (338,347235 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,19
12,533024 рдПрдордПрд╕ (334,660176 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,22
12,703328 рдПрдордПрд╕ (330,173626 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,25
12,884512 рдПрдордПрд╕ (325,530693 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,28
13,033472 рдПрдордПрд╕ (321,810182 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,31
13,239296 рдПрдордПрд╕ (316,807174 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,34
13,392448 рдПрдордПрд╕ (313,184256 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,37
13,624000 рдПрдордПрд╕ (307,861434 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,41
13,875520 рдПрдордПрд╕ (302,280855 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,44
14,126528 рдПрдордПрд╕ (296,909756 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,47
14,399328 рдПрдордПрд╕ (291,284699 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,50
14,690304 рдПрдордПрд╕ (285,515123 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,53
15,039136 рдПрдордПрд╕ (278,892623 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,56
15,478656 рдПрдордПрд╕ (270,973402 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,59
15,985664 рдПрдордПрд╕ (262,379092 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,62
16,668673 рдПрдордПрд╕ (251,627968 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,66
17,587200 рдПрдордПрд╕ (238,486174 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,69
18,690048 рдПрдордПрд╕ (224,413765 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,72
20,278816 рдПрдордПрд╕ (206,831789 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,75
22,545408 рдПрдордПрд╕ (186,038058 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,78
26,053312 рдПрдордПрд╕ (160,989275 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,81
31,895008 рдПрдордПрд╕ (131,503463 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,84
42,103294 рдПрдордПрд╕ (99,619378 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,87
61,849056 рдПрдордПрд╕ (67,815164 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,90
105,695999 рдПрдордПрд╕ (39,682713 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

0,94
240,204636 рдПрдордПрд╕ (17,461378 рдорд┐рд▓рд┐рдпрди рдХреБрдВрдЬреА/рд╕реЗрдХрдВрдб)

рдЬреИрд╕реЗ-рдЬреИрд╕реЗ рдЙрдкрдпреЛрдЧрд┐рддрд╛ рдмрдврд╝рддреА рд╣реИ, рдкреНрд░рджрд░реНрд╢рди рдХрдо рд╣реЛрддрд╛ рдЬрд╛рддрд╛ рд╣реИред рдЕрдзрд┐рдХрд╛рдВрд╢ рдорд╛рдорд▓реЛрдВ рдореЗрдВ рдпрд╣ рд╡рд╛рдВрдЫрдиреАрдп рдирд╣реАрдВ рд╣реИ. рдпрджрд┐ рдХреЛрдИ рдПрдкреНрд▓рд┐рдХреЗрд╢рди рдХрд┐рд╕реА рддрд╛рд▓рд┐рдХрд╛ рдореЗрдВ рддрддреНрд╡реЛрдВ рдХреЛ рд╕рдореНрдорд┐рд▓рд┐рдд рдХрд░рддрд╛ рд╣реИ рдФрд░ рдлрд┐рд░ рдЙрдиреНрд╣реЗрдВ рд╣рдЯрд╛ рджреЗрддрд╛ рд╣реИ (рдЙрджрд╛рд╣рд░рдг рдХреЗ рд▓рд┐рдП, рдХрд┐рд╕реА рдкреБрд╕реНрддрдХ рдореЗрдВ рд╢рдмреНрджреЛрдВ рдХреА рдЧрд┐рдирддреА рдХрд░рддреЗ рд╕рдордп), рддреЛ рдпрд╣ рдХреЛрдИ рд╕рдорд╕реНрдпрд╛ рдирд╣реАрдВ рд╣реИред рд▓реЗрдХрд┐рди рдпрджрд┐ рдПрдкреНрд▓рд┐рдХреЗрд╢рди рд▓рдВрдмреЗ рд╕рдордп рддрдХ рдЪрд▓рдиреЗ рд╡рд╛рд▓реА рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдХрд╛ рдЙрдкрдпреЛрдЧ рдХрд░рддрд╛ рд╣реИ (рдЙрджрд╛рд╣рд░рдг рдХреЗ рд▓рд┐рдП, рдЧреНрд░рд╛рдлрд┐рдХреНрд╕ рд╕рдВрдкрд╛рджрдХ рдореЗрдВ рдЫрд╡рд┐рдпреЛрдВ рдХреЗ рдЧреИрд░-рдЦрд╛рд▓реА рд╣рд┐рд╕реНрд╕реЛрдВ рдХреЛ рд╕рдВрдЧреНрд░рд╣реАрдд рдХрд░рдиреЗ рдХреЗ рд▓рд┐рдП рдЬрд╣рд╛рдВ рдЙрдкрдпреЛрдЧрдХрд░реНрддрд╛ рдЕрдХреНрд╕рд░ рдЬрд╛рдирдХрд╛рд░реА рдбрд╛рд▓рддрд╛ рд╣реИ рдФрд░ рд╣рдЯрд╛ рджреЗрддрд╛ рд╣реИ), рддреЛ рдпрд╣ рд╡реНрдпрд╡рд╣рд╛рд░ рд╕рдорд╕реНрдпрд╛рдЧреНрд░рд╕реНрдд рд╣реЛ рд╕рдХрддрд╛ рд╣реИред

рдФрд░ 64 рдорд┐рд▓рд┐рдпрди рдЗрдВрд╕рд░реНрдЯ (рдЙрдкрдпреЛрдЧ рдХрд╛рд░рдХ 0,5) рдХреЗ рдмрд╛рдж рдЧрд╣рд░рд╛рдИ рдХреА рдЬрд╛рдВрдЪ рдХрд░рдиреЗ рд╡рд╛рд▓реА рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдХреЛ рдорд╛рдкрд╛ред рдФрд╕рдд рдЧрд╣рд░рд╛рдИ 0,4774 рдереА, рдЗрд╕рд▓рд┐рдП рдЕрдзрд┐рдХрд╛рдВрд╢ рдЪрд╛рдмрд┐рдпрд╛рдБ рдпрд╛ рддреЛ рд╕рд░реНрд╡реЛрддреНрддрдо рд╕рдВрднрд╡ рд╕реНрд▓реЙрдЯ рдореЗрдВ рдереАрдВ рдпрд╛ рд╕рд░реНрд╡реЛрддреНрддрдо рд╕реНрдерд┐рддрд┐ рд╕реЗ рдПрдХ рд╕реНрд▓реЙрдЯ рджреВрд░ рдереАрдВред рдЕрдзрд┐рдХрддрдо рдзреНрд╡рдирд┐ рдЧрд╣рд░рд╛рдИ 60 рдереА.

рдлрд┐рд░ рдореИрдВрдиреЗ 124 рдорд┐рд▓рд┐рдпрди рдЗрдВрд╕рд░реНрдЯ (рдЙрдкрдпреЛрдЧ рдХрд╛рд░рдХ 0,97) рдХреЗ рд╕рд╛рде рдПрдХ рдореЗрдЬ рдкрд░ рдЬрд╛рдВрдЪ рдХреА рдЧрд╣рд░рд╛рдИ рдХреЛ рдорд╛рдкрд╛ред рдФрд╕рдд рдЧрд╣рд░рд╛рдИ рдкрд╣рд▓реЗ рд╕реЗ рд╣реА 10,1757 рдереА, рдФрд░ рдЕрдзрд┐рдХрддрдо - 6474 (!!). рдЙрдЪреНрдЪ рдЙрдкрдпреЛрдЧ рджрд░ рдкрд░ рд░реИрдЦрд┐рдХ рд╕рдВрд╡реЗрджрди рдкреНрд░рджрд░реНрд╢рди рдореЗрдВ рдХрд╛рдлреА рдЧрд┐рд░рд╛рд╡рдЯ рдЖрддреА рд╣реИред

рдЗрд╕ рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдХреА рдЙрдкрдпреЛрдЧ рджрд░ рдХреЛ рдХрдо рд░рдЦрдирд╛ рд╕рдмрд╕реЗ рдЕрдЪреНрдЫрд╛ рд╣реИред рд▓реЗрдХрд┐рди рдлрд┐рд░ рд╣рдо рдореЗрдореЛрд░реА рдЦрдкрдд рдХреА рдХреАрдордд рдкрд░ рдкреНрд░рджрд░реНрд╢рди рдмрдврд╝рд╛рддреЗ рд╣реИрдВред рд╕реМрднрд╛рдЧреНрдп рд╕реЗ, 32-рдмрд┐рдЯ рдХреБрдВрдЬрд┐рдпреЛрдВ рдФрд░ рдорд╛рдиреЛрдВ рдХреЗ рдорд╛рдорд▓реЗ рдореЗрдВ, рдЗрд╕реЗ рдЙрдЪрд┐рдд рдард╣рд░рд╛рдпрд╛ рдЬрд╛ рд╕рдХрддрд╛ рд╣реИред рдпрджрд┐ рдЙрдкрд░реЛрдХреНрдд рдЙрджрд╛рд╣рд░рдг рдореЗрдВ, 128 рдорд┐рд▓рд┐рдпрди рддрддреНрд╡реЛрдВ рд╡рд╛рд▓реА рддрд╛рд▓рд┐рдХрд╛ рдореЗрдВ, рд╣рдо рдЙрдкрдпреЛрдЧ рдХрд╛рд░рдХ 0,25 рд░рдЦрддреЗ рд╣реИрдВ, рддреЛ рд╣рдо рдЗрд╕рдореЗрдВ 32 рдорд┐рд▓рд┐рдпрди рд╕реЗ рдЕрдзрд┐рдХ рддрддреНрд╡ рдирд╣реАрдВ рд░рдЦ рд╕рдХрддреЗ рд╣реИрдВ, рдФрд░ рд╢реЗрд╖ 96 рдорд┐рд▓рд┐рдпрди рд╕реНрд▓реЙрдЯ рдЦреЛ рдЬрд╛рдПрдВрдЧреЗ - рдкреНрд░рддреНрдпреЗрдХ рдЬреЛрдбрд╝реА рдХреЗ рд▓рд┐рдП 8 рдмрд╛рдЗрдЯреНрд╕ , 768 рдПрдордмреА рдХреА рдЦреЛрдИ рд╣реБрдИ рдореЗрдореЛрд░реАред

рдХреГрдкрдпрд╛ рдзреНрдпрд╛рди рджреЗрдВ рдХрд┐ рд╣рдо рд╡реАрдбрд┐рдпреЛ рдХрд╛рд░реНрдб рдореЗрдореЛрд░реА рдХреЗ рдиреБрдХрд╕рд╛рди рдХреЗ рдмрд╛рд░реЗ рдореЗрдВ рдмрд╛рдд рдХрд░ рд░рд╣реЗ рд╣реИрдВ, рдЬреЛ рд╕рд┐рд╕реНрдЯрдо рдореЗрдореЛрд░реА рд╕реЗ рдЕрдзрд┐рдХ рдореВрд▓реНрдпрд╡рд╛рди рд╕рдВрд╕рд╛рдзрди рд╣реИред рд╣рд╛рд▓рд╛рдБрдХрд┐ CUDA рдХрд╛ рд╕рдорд░реНрдерди рдХрд░рдиреЗ рд╡рд╛рд▓реЗ рдЕрдзрд┐рдХрд╛рдВрд╢ рдЖрдзреБрдирд┐рдХ рдбреЗрд╕реНрдХрдЯреЙрдк рдЧреНрд░рд╛рдлрд╝рд┐рдХреНрд╕ рдХрд╛рд░реНрдб рдореЗрдВ рдХрдо рд╕реЗ рдХрдо 4 рдЬреАрдмреА рдореЗрдореЛрд░реА рд╣реЛрддреА рд╣реИ (рд▓реЗрдЦрди рдХреЗ рд╕рдордп, NVIDIA 2080 Ti рдореЗрдВ 11 рдЬреАрдмреА рд╣реИ), рдлрд┐рд░ рднреА рдЗрддрдиреА рдорд╛рддреНрд░рд╛ рдЦреЛрдирд╛ рд╕рдмрд╕реЗ рдмреБрджреНрдзрд┐рдорд╛рдиреА рднрд░рд╛ рдирд┐рд░реНрдгрдп рдирд╣реАрдВ рд╣реЛрдЧрд╛ред

рдмрд╛рдж рдореЗрдВ рдореИрдВ рдЙрди рд╡реАрдбрд┐рдпреЛ рдХрд╛рд░реНрдбреЛрдВ рдХреЗ рд▓рд┐рдП рд╣реИрд╢ рдЯреЗрдмрд▓ рдмрдирд╛рдиреЗ рдХреЗ рдмрд╛рд░реЗ рдореЗрдВ рдФрд░ рдЕрдзрд┐рдХ рд▓рд┐рдЦреВрдВрдЧрд╛ рдЬрд┐рдирдореЗрдВ рдЬрд╛рдВрдЪ рдХреА рдЧрд╣рд░рд╛рдИ рдХреА рд╕рдорд╕реНрдпрд╛ рдирд╣реАрдВ рд╣реИ, рд╕рд╛рде рд╣реА рд╣рдЯрд╛рдП рдЧрдП рд╕реНрд▓реЙрдЯ рдХрд╛ рдкреБрди: рдЙрдкрдпреЛрдЧ рдХрд░рдиреЗ рдХреЗ рддрд░реАрдХреЛрдВ рдХреЗ рдмрд╛рд░реЗ рдореЗрдВ рднреА рд▓рд┐рдЦреВрдВрдЧрд╛ред

рдзреНрд╡рдирд┐ рдЧрд╣рд░рд╛рдИ рдорд╛рдк

рдХрд┐рд╕реА рдХреБрдВрдЬреА рдХреА рдЬрд╛рдВрдЪ рдЧрд╣рд░рд╛рдИ рдирд┐рд░реНрдзрд╛рд░рд┐рдд рдХрд░рдиреЗ рдХреЗ рд▓рд┐рдП, рд╣рдо рдХреБрдВрдЬреА рдХреЗ рд╣реИрд╢ (рдЗрд╕рдХрд╛ рдЖрджрд░реНрд╢ рддрд╛рд▓рд┐рдХрд╛ рд╕реВрдЪрдХрд╛рдВрдХ) рдХреЛ рдЙрд╕рдХреЗ рд╡рд╛рд╕реНрддрд╡рд┐рдХ рддрд╛рд▓рд┐рдХрд╛ рд╕реВрдЪрдХрд╛рдВрдХ рд╕реЗ рдирд┐рдХрд╛рд▓ рд╕рдХрддреЗ рд╣реИрдВ:

// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);

рджреЛ рдХреЗ рджреЛ рдХреЗ рдкреВрд░рдХ рдмрд╛рдЗрдирд░реА рд╕рдВрдЦреНрдпрд╛рдУрдВ рдХреЗ рдЬрд╛рджреВ рдФрд░ рдЗрд╕ рддрдереНрдп рдХреЗ рдХрд╛рд░рдг рдХрд┐ рд╣реИрд╢ рддрд╛рд▓рд┐рдХрд╛ рдХреА рдХреНрд╖рдорддрд╛ рджреЛ рдХреА рд╢рдХреНрддрд┐ рджреЛ рд╣реИ, рдпрд╣ рджреГрд╖реНрдЯрд┐рдХреЛрдг рддрдм рднреА рдХрд╛рдо рдХрд░реЗрдЧрд╛ рдЬрдм рдХреБрдВрдЬреА рд╕реВрдЪрдХрд╛рдВрдХ рдХреЛ рддрд╛рд▓рд┐рдХрд╛ рдХреА рд╢реБрд░реБрдЖрдд рдореЗрдВ рд▓реЗ рдЬрд╛рдпрд╛ рдЬрд╛рдПрдЧрд╛ред рдЖрдЗрдП рдПрдХ рдХреБрдВрдЬреА рд▓реЗрдВ рдЬреЛ 1 рдкрд░ рд╣реИрд╢ рдХреА рдЧрдИ рд╣реИ, рд▓реЗрдХрд┐рди рд╕реНрд▓реЙрдЯ 3 рдореЗрдВ рдбрд╛рд▓реА рдЧрдИ рд╣реИред рдлрд┐рд░ рдХреНрд╖рдорддрд╛ 4 рд╡рд╛рд▓реА рдПрдХ рдЯреЗрдмрд▓ рдХреЗ рд▓рд┐рдП рд╣рдореЗрдВ рдорд┐рд▓рддрд╛ рд╣реИ (3 тАФ 1) & 3, рдЬреЛ 2 рдХреЗ рдмрд░рд╛рдмрд░ рд╣реИред

рдирд┐рд╖реНрдХрд░реНрд╖

рдпрджрд┐ рдЖрдкрдХреЗ рдХреЛрдИ рдкреНрд░рд╢реНрди рдпрд╛ рдЯрд┐рдкреНрдкрдгрд┐рдпрд╛рдБ рд╣реИрдВ, рддреЛ рдХреГрдкрдпрд╛ рдореБрдЭреЗ рдИрдореЗрд▓ рдХрд░реЗрдВ рдЯреНрд╡рд┐рдЯрд░ рдпрд╛ рдХреЛрдИ рдирдпрд╛ рд╡рд┐рд╖рдп рдЦреЛрд▓реЗрдВ рдЦрдЬрд╛рдиреЗ.

рдпрд╣ рдХреЛрдб рдЙрддреНрдХреГрд╖реНрдЯ рд▓реЗрдЦреЛрдВ рд╕реЗ рдкреНрд░реЗрд░рдгрд╛ рд▓реЗрдХрд░ рд▓рд┐рдЦрд╛ рдЧрдпрд╛ рдерд╛:

рднрд╡рд┐рд╖реНрдп рдореЗрдВ, рдореИрдВ рд╡реАрдбрд┐рдпреЛ рдХрд╛рд░реНрдб рдХреЗ рд▓рд┐рдП рд╣реИрд╢ рдЯреЗрдмрд▓ рдХрд╛рд░реНрдпрд╛рдиреНрд╡рдпрди рдХреЗ рдмрд╛рд░реЗ рдореЗрдВ рд▓рд┐рдЦрдирд╛ рдФрд░ рдЙрдирдХреЗ рдкреНрд░рджрд░реНрд╢рди рдХрд╛ рд╡рд┐рд╢реНрд▓реЗрд╖рдг рдХрд░рдирд╛ рдЬрд╛рд░реА рд░рдЦреВрдВрдЧрд╛ред рдореЗрд░реА рдпреЛрдЬрдирд╛рдУрдВ рдореЗрдВ рдбреЗрдЯрд╛ рд╕рдВрд░рдЪрдирд╛рдУрдВ рдореЗрдВ рдкрд░рдорд╛рдгреБ рд╕рдВрдЪрд╛рд▓рди рдХрд╛ рдЙрдкрдпреЛрдЧ рдХрд░рдХреЗ рдЪреЗрдирд┐рдВрдЧ, рд░реЙрдмрд┐рди рд╣реБрдб рд╣реИрд╢рд┐рдВрдЧ рдФрд░ рдХреБрдХреНрдХреВ рд╣реИрд╢рд┐рдВрдЧ рд╢рд╛рдорд┐рд▓ рд╣реИ рдЬреЛ рдЬреАрдкреАрдпреВ рдХреЗ рдЕрдиреБрдХреВрд▓ рд╣реИрдВред

рд╕реНрд░реЛрдд: www.habr.com

рдПрдХ рдЯрд┐рдкреНрдкрдгреА рдЬреЛрдбрд╝реЗрдВ