GPU рд╕рд╛рдареА рд╕рд╛рдзреЗ рд╣реЕрд╢ рдЯреЗрдмрд▓

GPU рд╕рд╛рдареА рд╕рд╛рдзреЗ рд╣реЕрд╢ рдЯреЗрдмрд▓
рдореА рддреЗ Github рд╡рд░ рдкреЛрд╕реНрдЯ рдХреЗрд▓реЗ рдЖрд╣реЗ рдирд╡реАрди рдкреНрд░рдХрд▓реНрдк рдПрдХ рд╕рд╛рдзрд╛ GPU рд╣реЕрд╢ рдЯреЗрдмрд▓.

рд╣реЗ рдПрдХ рд╕рд╛рдзреЗ GPU рд╣реЕрд╢ рдЯреЗрдмрд▓ рдЖрд╣реЗ рдЬреЗ рдкреНрд░рддрд┐ рд╕реЗрдХрдВрдж рд╢реЗрдХрдбреЛ рд▓рд╛рдЦреЛ рдЗрдиреНрд╕рд░реНрдЯрд╡рд░ рдкреНрд░рдХреНрд░рд┐рдпрд╛ рдХрд░рдгреНрдпрд╛рд╕ рд╕рдХреНрд╖рдо рдЖрд╣реЗ. рдорд╛рдЭреНрдпрд╛ NVIDIA GTX 1060 рд▓реЕрдкрдЯреЙрдкрд╡рд░, рдХреЛрдб 64 рджрд╢рд▓рдХреНрд╖ рдпрд╛рджреГрдЪреНрдЫрд┐рдХрдкрдгреЗ рд╡реНрдпреБрддреНрдкрдиреНрди рдХреЗрд▓реЗрд▓реНрдпрд╛ рдХреА-рд╡реНрд╣реЕрд▓реНрдпреВ рдЬреЛрдбреНрдпрд╛ рд╕реБрдорд╛рд░реЗ 210 ms рдордзреНрдпреЗ рдШрд╛рд▓рддреЛ рдЖрдгрд┐ рд╕реБрдорд╛рд░реЗ 32 ms рдордзреНрдпреЗ 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() рд╣реЕрд╢ рдЯреЗрдмрд▓рдордзреАрд▓ рдЬреЛрдбреНрдпрд╛рдВрдЪреНрдпрд╛ рдЕреЕрд░реЗрд╕рд╛рдареА рд╕рд╛рдзрд╛ рдХрд┐рдВрд╡рд╛ рдЕрд╕реНрдерд┐рд░ рдкреЙрдЗрдВрдЯрд░ рд╡рд╛рдкрд░рд╛. CUDA рджрд╕реНрддрдРрд╡рдЬреАрдХрд░рдг рдЕрд╕реЗ рд╕рд╛рдВрдЧрддреЗ рдХреА:

рдХрдВрдкрд╛рдЗрд▓рд░ рдЬрд╛рдЧрддрд┐рдХ рдХрд┐рдВрд╡рд╛ рд╕рд╛рдорд╛рдпрд┐рдХ рдореЗрдорд░реАрдордзреНрдпреЗ рд╡рд╛рдЪрди рдЖрдгрд┐ рд▓реЗрдЦрди рдСрдкреНрдЯрд┐рдорд╛рдЗрдЭ рдХрд░рдгреЗ рдирд┐рд╡рдбреВ рд╢рдХрддреЛ... рд╣реЗ рдСрдкреНрдЯрд┐рдорд╛рдпрдЭреЗрд╢рди рдХреАрд╡рд░реНрдб рд╡рд╛рдкрд░реВрди рдЕрдХреНрд╖рдо рдХреЗрд▓реЗ рдЬрд╛рдК рд╢рдХрддрд╛рдд volatile: ... рдпрд╛ рд╡реНрд╣реЗрд░рд┐рдПрдмрд▓рдЪрд╛ рдХреЛрдгрддрд╛рд╣реА рд╕рдВрджрд░реНрдн рд░рд┐рдЕрд▓ рдореЗрдорд░реА рд░реАрдб рдХрд┐рдВрд╡рд╛ рд░рд╛рдЗрдЯ рдЗрдВрд╕реНрдЯреНрд░рдХреНрд╢рдирдордзреНрдпреЗ рд╕рдВрдХрд▓рд┐рдд рдХреЗрд▓рд╛ рдЬрд╛рддреЛ.

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

рдЙрддреНрдкрд╛рджрдХрддрд╛

64 рджрд╢рд▓рдХреНрд╖ рдШрдЯрдХ рдШрд╛рд▓рдгреНрдпрд╛рд╕рд╛рдареА рдЖрдгрд┐ рддреНрдпрд╛рдкреИрдХреА 32 рджрд╢рд▓рдХреНрд╖ рд╣рдЯрд╡рд┐рдгреНрдпрд╛рдЪреНрдпрд╛ рдЪрд╛рдЪрдгреАрдордзреНрдпреЗ, рджрд░рдореНрдпрд╛рди рд╕реНрдкрд░реНрдзрд╛ std::unordered_map рдЖрдгрд┐ GPU рд╕рд╛рдареА рдЕрдХреНрд╖рд░рд╢рдГ рд╣реЕрд╢ рдЯреЗрдмрд▓ рдирд╛рд╣реА:

GPU рд╕рд╛рдареА рд╕рд╛рдзреЗ рд╣реЕрд╢ рдЯреЗрдмрд▓
std::unordered_map рдШрдЯрдХ рдШрд╛рд▓рдгреНрдпрд╛рд╕рд╛рдареА рдЖрдгрд┐ рдХрд╛рдвреВрди рдЯрд╛рдХрдгреНрдпрд╛рд╕рд╛рдареА рдЖрдгрд┐ рдирдВрддрд░ рддреНрдпрд╛рдВрдирд╛ рдореБрдХреНрдд рдХрд░рдгреНрдпрд╛рд╕рд╛рдареА 70 ms рдЦрд░реНрдЪ рдХреЗрд▓реЗ unordered_map (рд▓рд╛рдЦреЛ рдШрдЯрдХрд╛рдВрдкрд╛рд╕реВрди рдореБрдХреНрдд рд╣реЛрдгреНрдпрд╛рд╕рд╛рдареА рдЦреВрдк рд╡реЗрд│ рд▓рд╛рдЧрддреЛ, рдХрд╛рд░рдг рдЖрдд unordered_map рдПрдХрд╛рдзрд┐рдХ рдореЗрдорд░реА рд╡рд╛рдЯрдк рдХреЗрд▓реЗ рдЬрд╛рддреЗ). рдкреНрд░рд╛рдорд╛рдгрд┐рдХрдкрдгреЗ рд╕рд╛рдВрдЧрд╛рдпрдЪреЗ рддрд░, std:unordered_map рдкреВрд░реНрдгрдкрдгреЗ рднрд┐рдиреНрди рдирд┐рд░реНрдмрдВрдз. рд╣рд╛ рдПрдХрдЪ CPU рдереНрд░реЗрдб рдЖрд╣реЗ, рдЬреЛ рдХреЛрдгрддреНрдпрд╛рд╣реА рдЖрдХрд╛рд░рд╛рдЪреНрдпрд╛ рдХреА-рд╡реНрд╣реЕрд▓реНрдпреВрд▓рд╛ рд╕рдкреЛрд░реНрдЯ рдХрд░рддреЛ, рдЙрдЪреНрдЪ рд╡рд╛рдкрд░ рджрд░рд╛рдВрд╡рд░ рдЪрд╛рдВрдЧрд▓реА рдХрд╛рдордЧрд┐рд░реА рдХрд░рддреЛ рдЖрдгрд┐ рдПрдХрд╛рдзрд┐рдХ рд╣рдЯрд╡рд▓реНрдпрд╛рдирдВрддрд░ рд╕реНрдерд┐рд░ рдХрд╛рдордЧрд┐рд░реА рджрд╛рдЦрд╡рддреЛ.

GPU рдЖрдгрд┐ рдЗрдВрдЯрд░-рдкреНрд░реЛрдЧреНрд░рд╛рдо рдХрдореНрдпреБрдирд┐рдХреЗрд╢рдирд╕рд╛рдареА рд╣реЕрд╢ рдЯреЗрдмрд▓рдЪрд╛ рдХрд╛рд▓рд╛рд╡рдзреА 984 ms рд╣реЛрддрд╛. рдпрд╛рдд рдЯреЗрдмрд▓ рдореЗрдорд░реАрдордзреНрдпреЗ рдареЗрд╡рдгреНрдпрд╛рдд рдЖрдгрд┐ рд╣рдЯрд╡рдгреНрдпрд╛рдд рдШрд╛рд▓рд╡рд▓реЗрд▓рд╛ рд╡реЗрд│ (рдПрдХрджрд╛ 1 GB рдореЗрдорд░реА рд╡рд╛рдЯрдк рдХрд░рдгреЗ, рдЬреНрдпрд╛рд▓рд╛ CUDA рдордзреНрдпреЗ рдереЛрдбрд╛ рд╡реЗрд│ рд▓рд╛рдЧрддреЛ), рдШрдЯрдХ рдШрд╛рд▓рдгреЗ рдЖрдгрд┐ рд╣рдЯрд╡рдгреЗ рдЖрдгрд┐ рддреНрдпрд╛рд╡рд░ рдкреБрдирд░рд╛рд╡реГрддреНрддреА рдХрд░рдгреЗ рд╕рдорд╛рд╡рд┐рд╖реНрдЯ рдЖрд╣реЗ. рд╡реНрд╣рд┐рдбреАрдУ рдХрд╛рд░реНрдб рдореЗрдорд░реАрдордзреАрд▓ рдЖрдгрд┐ рддреНрдпрд╛рддреАрд▓ рд╕рд░реНрд╡ рдкреНрд░рддреА рджреЗрдЦреАрд▓ рд╡рд┐рдЪрд╛рд░рд╛рдд рдШреЗрддрд▓реНрдпрд╛ рдЬрд╛рддрд╛рдд.

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

рд╡реНрд╣рд┐рдбрд┐рдУ рдХрд╛рд░реНрдбрд╕рд╛рдареА рд╣реЕрд╢ рдЯреЗрдмрд▓ рдЙрдЪреНрдЪ рдереНрд░реВрдкреБрдЯ рдЖрдгрд┐ рд╕рдХреНрд░рд┐рдп рд╕рдорд╛рдВрддрд░реАрдХрд░рдгрд╛рдореБрд│реЗ рдЙрдЪреНрдЪ рдХрд╛рд░реНрдпрдХреНрд╖рдорддрд╛ рджрд░реНрд╢рд╡рддреЗ.

рдЙрдгреАрд╡рд╛

рд╣реЕрд╢ рдЯреЗрдмрд▓ рдЖрд░реНрдХрд┐рдЯреЗрдХреНрдЪрд░рдордзреНрдпреЗ рдХрд╛рд╣реА рд╕рдорд╕реНрдпрд╛ рдЖрд╣реЗрдд рдЬреНрдпрд╛рдВрдЪреА рдЬрд╛рдгреАрд╡ рдЕрд╕рдгреЗ рдЖрд╡рд╢реНрдпрдХ рдЖрд╣реЗ:

  • рдХреНрд▓рд╕реНрдЯрд░рд┐рдВрдЧрдореБрд│реЗ рд▓реАрдирд┐рдпрд░ рдкреНрд░реЛрдмрд┐рдВрдЧрд▓рд╛ рдЕрдбрдерд│рд╛ рдпреЗрддреЛ, рдЬреНрдпрд╛рдореБрд│реЗ рдЯреЗрдмрд▓рдордзреАрд▓ рдХрд│рд╛ рдЕрдЧрджреА рдХрдореА рдареЗрд╡рд▓реНрдпрд╛ рдЬрд╛рддрд╛рдд.
  • рдлрдВрдХреНрд╢рди рд╡рд╛рдкрд░реВрди рдХреА рдХрд╛рдврд▓реНрдпрд╛ рдЬрд╛рдд рдирд╛рд╣реАрдд delete рдЖрдгрд┐ рдХрд╛рд▓рд╛рдВрддрд░рд╛рдиреЗ рддреЗ рдЯреЗрдмрд▓ рдЧреЛрдВрдзрд│рддрд╛рдд.

рдкрд░рд┐рдгрд╛рдореА, рд╣реЕрд╢ рдЯреЗрдмрд▓рдЪреЗ рдХрд╛рд░реНрдпрдкреНрд░рджрд░реНрд╢рди рд╣рд│реВрд╣рд│реВ рдЦрд░рд╛рдм рд╣реЛрдК рд╢рдХрддреЗ, рд╡рд┐рд╢реЗрд╖рдд: рдЬрд░ рддреЗ рдмрд░реНрдпрд╛рдЪ рдХрд╛рд│рд╛рдкрд╛рд╕реВрди рдЕрд╕реНрддрд┐рддреНрд╡рд╛рдд рдЕрд╕реЗрд▓ рдЖрдгрд┐ рддреНрдпрд╛рдд рдЕрд╕рдВрдЦреНрдп рдЗрдиреНрд╕рд░реНрдЯ рдЖрдгрд┐ рдбрд┐рд▓реАрдЯ рдЕрд╕рддреАрд▓. рд╣реЗ рддреЛрдЯреЗ рдХрдореА рдХрд░рдгреНрдпрд╛рдЪрд╛ рдПрдХ рдорд╛рд░реНрдЧ рдореНрд╣рдгрдЬреЗ рдмрд░реНтАНрдпрд╛рдкреИрдХреА рдХрдореА рд╡рд╛рдкрд░ рджрд░рд╛рд╕рд╣ рдирд╡реАрди рдЯреЗрдмрд▓рдордзреНрдпреЗ рдкреБрдиреНрд╣рд╛ рд╣реЕрд╢ рдХрд░рдгреЗ рдЖрдгрд┐ рд░реАрд╣реЕрд╢рд┐рдВрдЧ рджрд░рдореНрдпрд╛рди рдХрд╛рдврд▓реЗрд▓реНрдпрд╛ рдХреА рдлрд┐рд▓реНрдЯрд░ рдХрд░рдгреЗ.

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

рд╡рд╛рдкрд░ рджрд░
рдЕрдВрддрд░реНрднреВрдд рдХрд╛рд▓рд╛рд╡рдзреА 4 рдШрдЯрдХ

0,00
11,608448 ms (361,314798 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,03
11,751424 ms (356,918799 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,06
11,942592 ms (351,205515 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,09
12,081120 ms (347,178429 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,12
12,242560 ms (342,600233 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,16
12,396448 ms (338,347235 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,19
12,533024 ms (334,660176 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,22
12,703328 ms (330,173626 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,25
12,884512 ms (325,530693 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,28
13,033472 ms (321,810182 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,31
13,239296 ms (316,807174 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,34
13,392448 ms (313,184256 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,37
13,624000 ms (307,861434 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,41
13,875520 ms (302,280855 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,44
14,126528 ms (296,909756 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,47
14,399328 ms (291,284699 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,50
14,690304 ms (285,515123 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,53
15,039136 ms (278,892623 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,56
15,478656 ms (270,973402 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,59
15,985664 ms (262,379092 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,62
16,668673 ms (251,627968 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,66
17,587200 ms (238,486174 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,69
18,690048 ms (224,413765 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,72
20,278816 ms (206,831789 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,75
22,545408 ms (186,038058 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,78
26,053312 ms (160,989275 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,81
31,895008 ms (131,503463 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,84
42,103294 ms (99,619378 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,87
61,849056 ms (67,815164 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,90
105,695999 ms (39,682713 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

0,94
240,204636 ms (17,461378 рджрд╢рд▓рдХреНрд╖ рдХреА/рд╕реЗ.)

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

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

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

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

рдХреГрдкрдпрд╛ рд▓рдХреНрд╖рд╛рдд рдШреНрдпрд╛ рдХреА рдЖрдореНрд╣реА рд╡реНрд╣рд┐рдбрд┐рдУ рдХрд╛рд░реНрдб рдореЗрдорд░реА рдЧрдорд╛рд╡рдгреНрдпрд╛рдмрджреНрджрд▓ рдмреЛрд▓рдд рдЖрд╣реЛрдд, рдЬреЗ рд╕рд┐рд╕реНрдЯрдо рдореЗрдорд░реАрдкреЗрдХреНрд╖рд╛ рдЕрдзрд┐рдХ рдореМрд▓реНрдпрд╡рд╛рди рд╕рдВрд╕рд╛рдзрди рдЖрд╣реЗ. рдЬрд░реА CUDA рдЪреЗ рд╕рдорд░реНрдерди рдХрд░рдгрд╛рд░реНтАНрдпрд╛ рдмрд╣реБрддреЗрдХ рдЖрдзреБрдирд┐рдХ рдбреЗрд╕реНрдХрдЯреЙрдк рдЧреНрд░рд╛рдлрд┐рдХреНрд╕ рдХрд╛рд░реНрдбреНрд╕рдордзреНрдпреЗ рдХрд┐рдорд╛рди 4 GB рдореЗрдорд░реА рдЕрд╕рддреЗ (рд▓реЗрдЦрдирд╛рдЪреНрдпрд╛ рд╡реЗрд│реА, NVIDIA 2080 Ti рдордзреНрдпреЗ 11 GB рдЕрд╕рддреЗ), рддрд░реАрд╣реА рдЕрд╢рд╛ рдкреНрд░рдорд╛рдгрд╛рдд рдЧрдорд╛рд╡рдгреЗ рд╣рд╛ рд╕рд░реНрд╡рд╛рдд рд╢рд╣рд╛рдгрд╛ рдирд┐рд░реНрдгрдп рдард░рдгрд╛рд░ рдирд╛рд╣реА.

рдирдВрддрд░ рдореА рд╡реНрд╣рд┐рдбрд┐рдУ рдХрд╛рд░реНрдбреНрд╕рд╕рд╛рдареА рд╣реЕрд╢ рдЯреЗрдмрд▓реНрд╕ рддрдпрд╛рд░ рдХрд░рдгреНрдпрд╛рдмрджреНрджрд▓ рдЕрдзрд┐рдХ рд▓рд┐рд╣реАрди рдЬреНрдпрд╛рдВрдирд╛ рдкреНрд░реЛрдмрд┐рдВрдЧ рдбреЗрдкреНрдердордзреНрдпреЗ рд╕рдорд╕реНрдпрд╛ рдпреЗрдд рдирд╛рд╣реАрдд, рддрд╕реЗрдЪ рд╣рдЯрд╡рд┐рд▓реЗрд▓реЗ рд╕реНрд▓реЙрдЯ рдкреБрдиреНрд╣рд╛ рд╡рд╛рдкрд░рдгреНрдпрд╛рдЪреЗ рдорд╛рд░реНрдЧ.

рдзреНрд╡рдиреА рдЦреЛрд▓реА рдореЛрдЬрдорд╛рдк

рдХреАрдЪреА рдкреНрд░реЛрдмрд┐рдВрдЧ рдбреЗрдкреНрде рдирд┐рд╢реНрдЪрд┐рдд рдХрд░рдгреНрдпрд╛рд╕рд╛рдареА, рдЖрдореНрд╣реА рддреНрдпрд╛рдЪреНрдпрд╛ рд╡рд╛рд╕реНрддрд╡рд┐рдХ рдЯреЗрдмрд▓ рдЗрдВрдбреЗрдХреНрд╕рдордзреВрди рдХреАрдЪрд╛ рд╣реЕрд╢ (рддреНрдпрд╛рдЪрд╛ рдЖрджрд░реНрд╢ рдЯреЗрдмрд▓ рдЗрдВрдбреЗрдХреНрд╕) рдХрд╛рдвреВ рд╢рдХрддреЛ:

// 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 рдЪреНрдпрд╛ рд╕рдорддреБрд▓реНрдп рдЖрд╣реЗ.

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

рдЖрдкрд▓реНрдпрд╛рдХрдбреЗ рдкреНрд░рд╢реНрди рдХрд┐рдВрд╡рд╛ рдЯрд┐рдкреНрдкрдгреНрдпрд╛ рдЕрд╕рд▓реНрдпрд╛рд╕, рдХреГрдкрдпрд╛ рдорд▓рд╛ рдпреЗрдереЗ рдИрдореЗрд▓ рдХрд░рд╛ Twitter рдХрд┐рдВрд╡рд╛ рдордзреНрдпреЗ рдПрдХ рдирд╡реАрди рд╡рд┐рд╖рдп рдЙрдШрдбрд╛ рднрд╛рдВрдбрд╛рд░.

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

рднрд╡рд┐рд╖реНрдпрд╛рдд, рдореА рд╡реНрд╣рд┐рдбрд┐рдУ рдХрд╛рд░реНрдбрд╕рд╛рдареА рд╣реЕрд╢ рдЯреЗрдмрд▓ рдЕрдВрдорд▓рдмрдЬрд╛рд╡рдгреАрдмрджреНрджрд▓ рд▓рд┐рд╣рд┐рдд рд░рд╛рд╣реАрди рдЖрдгрд┐ рддреНрдпрд╛рдВрдЪреНрдпрд╛ рдХрд╛рд░реНрдпрдкреНрд░рджрд░реНрд╢рдирд╛рдЪреЗ рд╡рд┐рд╢реНрд▓реЗрд╖рдг рдХрд░реЗрди. рдорд╛рдЭреНрдпрд╛ рдпреЛрдЬрдирд╛рдВрдордзреНрдпреЗ GPU рдЕрдиреБрдХреВрд▓ рдЕрд╕рд▓реЗрд▓реНрдпрд╛ рдбреЗрдЯрд╛ рд╕реНрдЯреНрд░рдХреНрдЪрд░реНрд╕рдордзреНрдпреЗ рдЕрдгреВ рдСрдкрд░реЗрд╢рдиреНрд╕ рд╡рд╛рдкрд░реВрди рдЪреЗрдирд┐рдВрдЧ, рд░реЙрдмрд┐рди рд╣реВрдб рд╣реЕрд╢рд┐рдВрдЧ рдЖрдгрд┐ рдХреБрдХреВ рд╣реЕрд╢рд┐рдВрдЧ рдпрд╛рдВрдЪрд╛ рд╕рдорд╛рд╡реЗрд╢ рдЖрд╣реЗ.

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

рдПрдХ рдЯрд┐рдкреНрдкрдгреА рдЬреЛрдбрд╛