GPU เจฒเจˆ เจธเจงเจพเจฐเจจ เจนเฉˆเจธเจผ เจŸเฉ‡เจฌเจฒ

GPU เจฒเจˆ เจธเจงเจพเจฐเจจ เจนเฉˆเจธเจผ เจŸเฉ‡เจฌเจฒ
เจฎเฉˆเจ‚ เจ‡เจธเจจเฉ‚เฉฐ Github 'เจคเฉ‡ เจชเฉ‹เจธเจŸ เจ•เฉ€เจคเจพ เจนเฉˆ เจจเจตเจพเจ‚ เจชเฉเจฐเฉ‹เจœเฉˆเจ•เจŸ เจ‡เฉฑเจ• เจธเจงเจพเจฐเจจ GPU เจนเฉˆเจธเจผ เจŸเฉ‡เจฌเจฒ.

เจ‡เจน เจ‡เฉฑเจ• เจธเจงเจพเจฐเจจ GPU เจนเฉˆเจธเจผ เจŸเฉ‡เจฌเจฒ เจนเฉˆ เจœเฉ‹ เจชเฉเจฐเจคเฉ€ เจธเจ•เจฟเฉฐเจŸ เจฒเฉฑเจ–เจพเจ‚ เจธเฉฐเจฎเจฟเจฒเจจเจพเจ‚ เจฆเฉ€ เจชเฉเจฐเจ•เจฟเจฐเจฟเจ† เจ•เจฐเจจ เจฆเฉ‡ เจธเจฎเจฐเฉฑเจฅ เจนเฉˆเฅค เจฎเฉ‡เจฐเฉ‡ NVIDIA GTX 1060 เจฒเฉˆเจชเจŸเจพเจช 'เจคเฉ‡, เจ•เฉ‹เจก เจฒเจ—เจญเจ— 64 ms เจตเจฟเฉฑเจš 210 เจฎเจฟเจฒเฉ€เจ…เจจ เจฌเฉ‡เจคเจฐเจคเฉ€เจฌเฉ‡ เจคเฉŒเจฐ 'เจคเฉ‡ เจคเจฟเจ†เจฐ เจ•เฉ€เจคเฉ‡ เจ•เฉเฉฐเจœเฉ€-เจฎเฉเฉฑเจฒ เจœเฉ‹เฉœเจฟเจ†เจ‚ เจจเฉ‚เฉฐ เจธเจผเจพเจฎเจฒ เจ•เจฐเจฆเจพ เจนเฉˆ เจ…เจคเฉ‡ เจฒเจ—เจญเจ— 32 ms เจตเจฟเฉฑเจš 64 เจฎเจฟเจฒเฉ€เจ…เจจ เจœเฉ‹เฉœเจฟเจ†เจ‚ เจจเฉ‚เฉฐ เจนเจŸเจพ เจฆเจฟเฉฐเจฆเจพ เจนเฉˆเฅค

เจญเจพเจต, เจ‡เฉฑเจ• เจฒเฉˆเจชเจŸเจพเจช 'เจคเฉ‡ เจธเจชเฉ€เจก เจฒเจ—เจญเจ— 300 เจฎเจฟเจฒเฉ€เจ…เจจ เจ‡เจจเจธเจฐเจŸเจธ/เจธเฉˆเจ•เฉฐเจก เจ…เจคเฉ‡ 500 เจฎเจฟเจฒเฉ€เจ…เจจ เจกเจฟเจฒเฉ€เจŸ/เจธเฉˆเจ•เฉฐเจก เจนเฉˆเฅค

เจธเจพเจฐเจฃเฉ€ เจจเฉ‚เฉฐ CUDA เจตเจฟเฉฑเจš เจฒเจฟเจ–เจฟเจ† เจ—เจฟเจ† เจนเฉˆ, เจนเจพเจฒเจพเจ‚เจ•เจฟ เจ‡เจนเฉ€ เจคเจ•เจจเฉ€เจ• HLSL เจœเจพเจ‚ GLSL 'เจคเฉ‡ เจฒเจพเจ—เฉ‚ เจ•เฉ€เจคเฉ€ เจœเจพ เจธเจ•เจฆเฉ€ เจนเฉˆเฅค เจตเฉ€เจกเฉ€เจ“ เจ•เจพเจฐเจก 'เจคเฉ‡ เจ‰เฉฑเจš เจชเฉเจฐเจฆเจฐเจธเจผเจจ เจจเฉ‚เฉฐ เจฏเจ•เฉ€เจจเฉ€ เจฌเจฃเจพเจ‰เจฃ เจฒเจˆ เจฒเจพเจ—เฉ‚ เจ•เจฐเจจ เจฆเฉ€เจ†เจ‚ เจ•เจˆ เจธเฉ€เจฎเจพเจตเจพเจ‚ เจนเจจ:

  • เจธเจฟเจฐเจซเจผ 32-เจฌเจฟเฉฑเจŸ เจ•เฉเฉฐเจœเฉ€เจ†เจ‚ เจ…เจคเฉ‡ เจธเจฎเจพเจจ เจฎเฉเฉฑเจฒเจพเจ‚ 'เจคเฉ‡ เจ•เจพเจฐเจตเจพเจˆ เจ•เฉ€เจคเฉ€ เจœเจพเจ‚เจฆเฉ€ เจนเฉˆเฅค
  • เจนเฉˆเจธเจผ เจŸเฉ‡เจฌเจฒ เจฆเจพ เจ‡เฉฑเจ• เจจเจฟเจธเจผเจšเจฟเจค เจ†เจ•เจพเจฐ เจนเฉˆเฅค
  • เจ…เจคเฉ‡ เจ‡เจน เจ†เจ•เจพเจฐ เจธเจผเจ•เจคเฉ€ เจฆเฉ‡ เจฆเฉ‹ เจฆเฉ‡ เจฌเจฐเจพเจฌเจฐ เจนเฉ‹เจฃเจพ เจšเจพเจนเฉ€เจฆเจพ เจนเฉˆ.

เจ•เฉเฉฐเจœเฉ€เจ†เจ‚ เจ…เจคเฉ‡ เจฎเฉเฉฑเจฒเจพเจ‚ เจฒเจˆ, เจคเฉเจนเจพเจจเฉ‚เฉฐ เจ‡เฉฑเจ• เจธเจงเจพเจฐเจจ เจกเฉˆเจฒเฉ€เจฎเฉ€เจŸเจฐ เจฎเจพเจฐเจ•เจฐ เจจเฉ‚เฉฐ เจฐเจฟเจœเจผเจฐเจต เจ•เจฐเจจ เจฆเฉ€ เจฒเฉ‹เฉœ เจนเฉˆ (เจ‰เจชเจฐเฉ‹เจ•เจค เจ•เฉ‹เจก เจตเจฟเฉฑเจš เจ‡เจน 0xffffffff เจนเฉˆ)เฅค

เจคเจพเจฒเฉ‡ เจคเฉ‹เจ‚ เจฌเจฟเจจเจพเจ‚ เจนเฉˆเจธเจผ เจŸเฉ‡เจฌเจฒ

เจนเฉˆเจธเจผ เจŸเฉ‡เจฌเจฒ เจ“เจชเจจ เจเจกเจฐเฉˆเจธเจฟเฉฐเจ— เจฆเฉ‡ เจจเจพเจฒ เจตเจฐเจคเจฆเจพ เจนเฉˆ เจฐเฉ‡เจ–เจฟเจ• เจชเฉœเจคเจพเจฒ, เจฏเจพเจจเฉ€, เจ‡เจน เจธเจฟเจฐเจซเจผ เจ•เฉเฉฐเจœเฉ€-เจฎเฉเฉฑเจฒ เจฆเฉ‡ เจœเฉ‹เฉœเจฟเจ†เจ‚ เจฆเฉ€ เจ‡เฉฑเจ• เจเจฐเฉ‡ เจนเฉˆ เจœเฉ‹ เจฎเฉˆเจฎเฉ‹เจฐเฉ€ เจตเจฟเฉฑเจš เจธเจŸเฉ‹เจฐ เจ•เฉ€เจคเฉ€ เจœเจพเจ‚เจฆเฉ€ เจนเฉˆ เจ…เจคเฉ‡ เจตเจงเฉ€เจ† เจ•เฉˆเจธเจผ เจชเฉเจฐเจฆเจฐเจธเจผเจจ เจนเฉˆเฅค เจšเฉ‡เจจเจฟเฉฐเจ— เจฒเจˆ เจตเฉ€ เจ‡เจนเฉ€ เจจเจนเฉ€เจ‚ เจ•เจฟเจนเจพ เจœเจพ เจธเจ•เจฆเจพ, เจœเจฟเจธ เจตเจฟเฉฑเจš เจฒเจฟเฉฐเจ•เจก เจธเฉ‚เจšเฉ€ เจตเจฟเฉฑเจš เจ‡เฉฑเจ• เจชเฉเจ†เจ‡เฉฐเจŸเจฐ เจฆเฉ€ เจ–เฉ‹เจœ เจ•เจฐเจจเจพ เจธเจผเจพเจฎเจฒ เจนเฉเฉฐเจฆเจพ เจนเฉˆเฅค เจ‡เฉฑเจ• เจนเฉˆเจธเจผ เจŸเฉ‡เจฌเจฒ เจ‡เฉฑเจ• เจธเจงเจพเจฐเจจ เจเจฐเฉ‡ เจธเจŸเฉ‹เจฐ เจ•เจฐเจจ เจตเจพเจฒเฉ‡ เจคเฉฑเจค เจนเฉˆ KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

เจŸเฉ‡เจฌเจฒ เจฆเจพ เจ†เจ•เจพเจฐ เจฆเฉ‹ เจฆเฉ€ เจธเจผเจ•เจคเฉ€ เจนเฉˆ, เจจเจพ เจ•เจฟ เจ‡เฉฑเจ• เจชเฉเจฐเจฎเฉเฉฑเจ– เจธเฉฐเจ–เจฟเจ†, เจ•เจฟเจ‰เจ‚เจ•เจฟ เจ‡เฉฑเจ• เจคเฉ‡เจœเจผ เจนเจฆเจพเจ‡เจค pow2/AND เจฎเจพเจธเจ• เจจเฉ‚เฉฐ เจฒเจพเจ—เฉ‚ เจ•เจฐเจจ เจฒเจˆ เจ•เจพเจซเจผเฉ€ เจนเฉˆ, เจชเจฐ เจฎเจพเจกเจฟเจŠเจฒเจธ เจ†เจชเจฐเฉ‡เจŸเจฐ เจฌเจนเฉเจค เจนเฉŒเจฒเฉ€ เจนเฉˆเฅค เจ‡เจน เจฒเฉ€เจจเฉ€เจ…เจฐ เจชเฉเจฐเฉ‹เจฌเจฟเฉฐเจ— เจฆเฉ‡ เจฎเจพเจฎเจฒเฉ‡ เจตเจฟเฉฑเจš เจฎเจนเฉฑเจคเจตเจชเฉ‚เจฐเจจ เจนเฉˆ, เจ•เจฟเจ‰เจ‚เจ•เจฟ เจ‡เฉฑเจ• เจฒเฉ€เจจเฉ€เจ…เจฐ เจŸเฉ‡เจฌเจฒ เจฒเฉเฉฑเจ•เจ…เฉฑเจช เจตเจฟเฉฑเจš เจธเจฒเจพเจŸ เจธเฉ‚เจšเจ•เจพเจ‚เจ• เจจเฉ‚เฉฐ เจนเจฐเฉ‡เจ• เจธเจฒเจพเจŸ เจตเจฟเฉฑเจš เจฒเจชเฉ‡เจŸเจฟเจ† เจœเจพเจฃเจพ เจšเจพเจนเฉ€เจฆเจพ เจนเฉˆเฅค เจ…เจคเฉ‡ เจจเจคเฉ€เจœเฉ‡ เจตเจœเฉ‹เจ‚, เจ•เจพเจฐเจตเจพเจˆ เจฆเฉ€ เจฒเจพเจ—เจค เจจเฉ‚เฉฐ เจนเจฐเฉ‡เจ• เจธเจฒเจพเจŸ เจตเจฟเฉฑเจš เจฎเฉ‹เจกเจฟเจŠเจฒเฉ‹ เจœเฉ‹เฉœเจฟเจ† เจœเจพเจ‚เจฆเจพ เจนเฉˆเฅค

เจธเจพเจฐเจฃเฉ€ เจธเจฟเจฐเจซเจผ เจนเจฐเฉ‡เจ• เจคเฉฑเจค เจฒเจˆ เจ•เฉเฉฐเจœเฉ€ เจ…เจคเฉ‡ เจฎเฉเฉฑเจฒ เจจเฉ‚เฉฐ เจธเจŸเฉ‹เจฐ เจ•เจฐเจฆเฉ€ เจนเฉˆ, เจ•เฉเฉฐเจœเฉ€ เจฆเจพ เจนเฉˆเจธเจผ เจจเจนเฉ€เจ‚เฅค เจ•เจฟเจ‰เจ‚เจ•เจฟ เจธเจพเจฐเจฃเฉ€ เจธเจฟเจฐเจซเจผ 32-เจฌเจฟเฉฑเจŸ เจ•เฉเฉฐเจœเฉ€เจ†เจ‚ เจจเฉ‚เฉฐ เจธเจŸเฉ‹เจฐ เจ•เจฐเจฆเฉ€ เจนเฉˆ, เจนเฉˆเจธเจผ เจฆเฉ€ เจ—เจฃเจจเจพ เจฌเจนเฉเจค เจคเฉ‡เจœเจผเฉ€ เจจเจพเจฒ เจ•เฉ€เจคเฉ€ เจœเจพเจ‚เจฆเฉ€ เจนเฉˆเฅค เจ‰เจชเจฐเฉ‹เจ•เจค เจ•เฉ‹เจก Murmur3 เจนเฉˆเจธเจผ เจฆเฉ€ เจตเจฐเจคเฉ‹เจ‚ เจ•เจฐเจฆเจพ เจนเฉˆ, เจœเฉ‹ เจธเจฟเจฐเจซ เจ•เฉเจ เจธเจผเจฟเจซเจŸเจพเจ‚, XORs เจ…เจคเฉ‡ เจ—เฉเจฃเจพ เจ•เจฐเจฆเจพ เจนเฉˆเฅค

เจนเฉˆเจธเจผ เจŸเฉ‡เจฌเจฒ เจฒเจพเจ•เจฟเฉฐเจ— เจธเฉเจฐเฉฑเจ–เจฟเจ† เจคเจ•เจจเฉ€เจ•เจพเจ‚ เจฆเฉ€ เจตเจฐเจคเฉ‹เจ‚ เจ•เจฐเจฆเจพ เจนเฉˆ เจœเฉ‹ เจฎเฉˆเจฎเฉ‹เจฐเฉ€ เจ†เจฐเจกเจฐ เจคเฉ‹เจ‚ เจธเฉเจคเฉฐเจคเจฐ เจนเจจเฅค เจญเจพเจตเฉ‡เจ‚ เจ•เฉเจ เจฐเจพเจˆเจŸ เจ“เจชเจฐเฉ‡เจธเจผเจจ เจ…เจœเจฟเจนเฉ‡ เจนเฉ‹เจฐ เจ“เจชเจฐเฉ‡เจธเจผเจจเจพเจ‚ เจฆเฉ‡ เจ•เฉเจฐเจฎ เจตเจฟเฉฑเจš เจตเจฟเจ˜เจจ เจชเจพเจ‰เจ‚เจฆเฉ‡ เจนเจจ, เจนเฉˆเจธเจผ เจŸเฉ‡เจฌเจฒ เจ…เจœเฉ‡ เจตเฉ€ เจธเจนเฉ€ เจธเจฅเจฟเจคเฉ€ เจจเฉ‚เฉฐ เจฌเจฐเจ•เจฐเจพเจฐ เจฐเฉฑเจ–เฉ‡เจ—เจพเฅค เจ…เจธเฉ€เจ‚ เจนเฉ‡เจ เจพเจ‚ เจ‡เจธ เจฌเจพเจฐเฉ‡ เจ—เฉฑเจฒ เจ•เจฐเจพเจ‚เจ—เฉ‡เฅค เจ‡เจน เจคเจ•เจจเฉ€เจ• เจตเฉ€เจกเฉ€เจ“ เจ•เจพเจฐเจกเจพเจ‚ เจจเจพเจฒ เจตเจงเฉ€เจ† เจ•เฉฐเจฎ เจ•เจฐเจฆเฉ€ เจนเฉˆ เจœเฉ‹ เจ‡เฉฑเจ•เฉ‹ เจธเจฎเฉ‡เจ‚ เจนเจœเจผเจพเจฐเจพเจ‚ เจฅเจฐเจฟเฉฑเจกเจพเจ‚ เจจเฉ‚เฉฐ เจšเจฒเจพเจ‰เจ‚เจฆเฉ‡ เจนเจจเฅค

เจนเฉˆเจธเจผ เจŸเฉ‡เจฌเจฒ เจตเจฟเฉฑเจš เจ•เฉเฉฐเจœเฉ€เจ†เจ‚ เจ…เจคเฉ‡ เจฎเฉเฉฑเจฒเจพเจ‚ เจจเฉ‚เฉฐ เจ–เจพเจฒเฉ€ เจ•เจฐเจจ เจฒเจˆ เจธเจผเฉเจฐเฉ‚ เจ•เฉ€เจคเจพ เจœเจพเจ‚เจฆเจพ เจนเฉˆเฅค

เจ•เฉ‹เจก เจจเฉ‚เฉฐ 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 เจ•เฉ‹เจก เจฆเฉเจ†เจฐเจพ เจตเจฐเจคเจฟเจ† เจœเจพเจตเฉ‡เจ—เจพ เจจเจพ เจ•เจฟ เจ•เฉ‡เจ‚เจฆเจฐเฉ€ เจชเฉเจฐเฉ‹เจธเฉˆเจธเจฐ เจฆเฉเจ†เจฐเจพ), เจคเจพเจ‚ เจŸเฉˆเจธเจŸ เจฆเจพ เจจเจคเฉ€เจœเจพ เจขเฉเจ•เจตเจพเจ‚ เจนเฉˆเฅค

เจตเฉ€เจกเฉ€เจ“ เจ•เจพเจฐเจก เจฒเจˆ เจนเฉˆเจธเจผ เจŸเฉ‡เจฌเจฒ เจ‰เฉฑเจš เจฅเฉเจฐเจฐเฉ‚เจชเฉเจŸ เจ…เจคเฉ‡ เจ•เจฟเจฐเจฟเจ†เจธเจผเฉ€เจฒ เจธเจฎเจพเจจเจคเจพ เจฆเฉ‡ เจ•เจพเจฐเจจ เจ‰เฉฑเจš เจชเฉเจฐเจฆเจฐเจธเจผเจจ เจจเฉ‚เฉฐ เจฆเจฐเจธเจพเจ‰เจ‚เจฆเจพ เจนเฉˆเฅค

shortcomings

เจนเฉˆเจธเจผ เจŸเฉ‡เจฌเจฒ เจ†เจฐเจ•เฉ€เจŸเฉˆเจ•เจšเจฐ เจตเจฟเฉฑเจš เจธเฉเจšเฉ‡เจค เจนเฉ‹เจฃ เจฒเจˆ เจ•เฉเจ เจฎเฉเฉฑเจฆเฉ‡ เจนเจจ:

  • เจฒเฉ€เจจเฉ€เจ…เจฐ เจชเฉเจฐเฉ‹เจฌเจฟเฉฐเจ— เจจเฉ‚เฉฐ เจ•เจฒเฉฑเจธเจŸเจฐเจฟเฉฐเจ— เจฆเฉเจ†เจฐเจพ เจ…เฉœเจฟเฉฑเจ•เจพ เจฌเจฃเจพเจ‡เจ† เจœเจพเจ‚เจฆเจพ เจนเฉˆ, เจœเจฟเจธ เจ•เจพเจฐเจจ เจธเจพเจฐเจฃเฉ€ เจตเจฟเฉฑเจš เจ•เฉเฉฐเจœเฉ€เจ†เจ‚ เจชเฉ‚เจฐเฉ€ เจคเจฐเฉเจนเจพเจ‚ เจ˜เฉฑเจŸ เจฐเฉฑเจ–เฉ€เจ†เจ‚ เจœเจพเจ‚เจฆเฉ€เจ†เจ‚ เจนเจจเฅค
  • เจซเฉฐเจ•เจธเจผเจจ เจฆเฉ€ เจตเจฐเจคเฉ‹เจ‚ เจ•เจฐเจ•เฉ‡ เจ•เฉเฉฐเจœเฉ€เจ†เจ‚ เจจเฉ‚เฉฐ เจจเจนเฉ€เจ‚ เจนเจŸเจพเจ‡เจ† เจœเจพเจ‚เจฆเจพ เจนเฉˆ 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 เจฆเฉ‡ เจฌเจฐเจพเจฌเจฐ เจนเฉˆเฅค

เจธเจฟเฉฑเจŸเจพ

เจœเฉ‡ เจคเฉเจนเจพเจกเฉ‡ เจ•เฉ‹เจˆ เจธเจตเจพเจฒ เจœเจพเจ‚ เจŸเจฟเฉฑเจชเจฃเฉ€เจ†เจ‚ เจนเจจ, เจคเจพเจ‚ เจ•เจฟเจฐเจชเจพ เจ•เจฐเจ•เฉ‡ เจฎเฉˆเจจเฉ‚เฉฐ เจ‡เจธ 'เจคเฉ‡ เจˆเจฎเฉ‡เจฒ เจ•เจฐเฉ‹ เจŸเจตเจฟเฉฑเจŸเจฐ เจœเจพเจ‚ เจตเจฟเฉฑเจš เจ‡เฉฑเจ• เจจเจตเจพเจ‚ เจตเจฟเจธเจผเจพ เจ–เฉ‹เจฒเฉเจนเฉ‹ เจฐเจฟเจชเฉ‹เจœเจผเจŸเจฐเฉ€เจ†เจ‚.

เจ‡เจน เจ•เฉ‹เจก เจธเจผเจพเจจเจฆเจพเจฐ เจฒเฉ‡เจ–เจพเจ‚ เจคเฉ‹เจ‚ เจชเฉเจฐเฉ‡เจฐเจจเจพ เจฒเฉˆ เจ•เฉ‡ เจฒเจฟเจ–เจฟเจ† เจ—เจฟเจ† เจธเฉ€:

เจญเจตเจฟเฉฑเจ– เจตเจฟเฉฑเจš, เจฎเฉˆเจ‚ เจตเฉ€เจกเฉ€เจ“ เจ•เจพเจฐเจกเจพเจ‚ เจฒเจˆ เจนเฉˆเจธเจผ เจŸเฉ‡เจฌเจฒ เจฒเจพเจ—เฉ‚ เจ•เจฐเจจ เจฌเจพเจฐเฉ‡ เจฒเจฟเจ–เจฃเจพ เจœเจพเจฐเฉ€ เจฐเฉฑเจ–เจพเจ‚เจ—เจพ เจ…เจคเฉ‡ เจ‰เจนเจจเจพเจ‚ เจฆเฉ‡ เจชเฉเจฐเจฆเจฐเจธเจผเจจ เจฆเจพ เจตเจฟเจธเจผเจฒเฉ‡เจธเจผเจฃ เจ•เจฐเจพเจ‚เจ—เจพ. เจฎเฉ‡เจฐเฉ€เจ†เจ‚ เจฏเฉ‹เจœเจจเจพเจตเจพเจ‚ เจตเจฟเฉฑเจš เจธเจผเจพเจฎเจฒ เจนเจจ เจšเฉ‡เจจเจฟเฉฐเจ—, เจฐเฉŒเจฌเจฟเจจ เจนเฉเฉฑเจก เจนเฉˆเจธเจผเจฟเฉฐเจ—, เจ…เจคเฉ‡ เจกเจพเจŸเจพ เจธเจŸเฉเจฐเจ•เจšเจฐเจœเจผ เจตเจฟเฉฑเจš เจชเจฐเจฎเจพเจฃเฉ‚ เจ“เจชเจฐเฉ‡เจธเจผเจจเจพเจ‚ เจฆเฉ€ เจตเจฐเจคเฉ‹เจ‚ เจ•เจฐเจฆเฉ‡ เจนเฉ‹เจ เจ•เฉ‹เจ•เฉ‚ เจนเฉˆเจธเจผเจฟเฉฐเจ— เจœเฉ‹ GPU เจ…เจจเฉเจ•เฉ‚เจฒ เจนเจจเฅค

เจธเจฐเฉ‹เจค: www.habr.com

เจ‡เฉฑเจ• เจŸเจฟเฉฑเจชเจฃเฉ€ เจœเฉ‹เฉœเฉ‹