פרסמתי את זה ב-Github פרויקט חדש A Simple GPU Hash Table.

זוהי טבלת Hash פשוטה של ​​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 סיביות, ה-hash מחושב מהר מאוד. הקוד שלמעלה משתמש ב-Hash Murmur3, שמבצע רק כמה משמרות, XOR והכפלות.

טבלת הגיבוב משתמשת בטכניקות הגנת נעילה שאינן תלויות בסדר הזיכרון. גם אם פעולות כתיבה מסוימות משבשות את הסדר של פעולות אחרות כאלה, טבלת ה-hash עדיין תשמור על המצב הנכון. נדבר על זה להלן. הטכניקה עובדת מצוין עם כרטיסי מסך שמריצים אלפי שרשורים במקביל.

המפתחות והערכים בטבלת ה-hash מאותחלים לריק.

ניתן לשנות את הקוד כך שיטפל גם במפתחות ובערכים של 64 סיביות. מפתחות דורשים פעולות קריאה, כתיבה והשוואה-החלפה אטומית. וערכים דורשים פעולות קריאה וכתיבה אטומיות. למרבה המזל, ב-CUDA, פעולות קריאה-כתיבה עבור ערכי 32 ו-64 סיביות הן אטומיות כל עוד הן מיושרות באופן טבעי (ראה להלן). כאן), וכרטיסי מסך מודרניים תומכים בפעולות השוואה והחלפה אטומיות של 64 סיביות. כמובן, כאשר עוברים ל-64 סיביות, הביצועים ירדו מעט.

מצב טבלת Hash

לכל זוג מפתח-ערך בטבלת גיבוב יכול להיות אחד מארבעה מצבים:

• מפתח וערך ריקים. במצב זה, טבלת ה-hash מאותחלת.
• המפתח נכתב, אך הערך עדיין לא נכתב. אם שרשור אחר קורא נתונים כעת, הוא יחזור ריק. זה נורמלי, אותו דבר היה קורה אם שרשור אחר של ביצוע היה עובד קצת קודם, ואנחנו מדברים על מבנה נתונים במקביל.
• גם המפתח וגם הערך נרשמים.
• הערך זמין לשרשורים אחרים של ביצוע, אך המפתח עדיין לא. זה יכול לקרות מכיוון שלמודל התכנות CUDA יש מודל זיכרון מסודר באופן רופף. זה נורמלי; בכל מקרה, המפתח עדיין ריק, גם אם הערך כבר לא כך.

ניואנס חשוב הוא שברגע שהמפתח נכתב לחריץ, הוא כבר לא זז - גם אם המפתח נמחק, נדבר על כך בהמשך.

קוד טבלת ה-hash אפילו עובד עם דגמי זיכרון מסודרים באופן רופף, שבהם סדר הקריאה והכתיבה של הזיכרון אינו ידוע. כאשר אנו מסתכלים על הכנסה, חיפוש ומחיקה בטבלת hash, זכור שכל זוג מפתח-ערך נמצא באחד מארבעת המצבים שתוארו לעיל.

הכנסה לטבלת גיבוב

פונקציית CUDA שמכניסה צמדי מפתח-ערך לטבלת hash נראית כך:

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);
    }
}

כדי להכניס מפתח, הקוד חוזר דרך מערך טבלת ה-hash החל מה-hash של המפתח שהוכנס. כל חריץ במערך מבצע פעולת השוואה והחלפה אטומית המשווה את המפתח בחריץ זה לריק. אם מזוהה אי התאמה, המפתח בחריץ מתעדכן עם המפתח שהוכנס, ולאחר מכן מוחזר מפתח החריץ המקורי. אם מפתח מקורי זה היה ריק או תואם למפתח שהוכנס, אזי הקוד מצא חריץ מתאים להכנסה והכניס את הערך שהוכנס לחריץ.

אם בקריאת קרנל אחת gpu_hashtable_insert() ישנם מספר אלמנטים עם אותו מפתח, ואז כל אחד מהערכים שלהם יכול להיכתב לחריץ המפתח. זה נחשב נורמלי: אחת מכתיבת ערך-מפתח במהלך הקריאה תצליח, אך מכיוון שכל זה קורה במקביל בתוך מספר שרשורים של ביצוע, איננו יכולים לחזות איזו כתיבת זיכרון תהיה האחרונה.

חיפוש טבלת Hash

קוד למפתחות חיפוש:

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);
        }
}

כדי למצוא את הערך של מפתח המאוחסן בטבלה, אנו עוברים דרך המערך החל מה-hash של המפתח שאנו מחפשים. בכל משבצת אנו בודקים אם המפתח הוא זה שאנו מחפשים, ואם כן, נחזיר את ערכו. אנו גם בודקים אם המפתח ריק, ואם כן, אנו מבטלים את החיפוש.

אם לא נוכל למצוא את המפתח, הקוד מחזיר ערך ריק.

ניתן לבצע את כל פעולות החיפוש הללו במקביל באמצעות הוספות ומחיקות. לכל זוג בטבלה יהיה אחד מארבעת המצבים שתוארו לעיל עבור הזרימה.

מחיקה בטבלת גיבוב

קוד למחיקת מפתחות:

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(), אלא שכאשר נמצא התאמה על מפתח, הוא הופך את הערך שלו לריק.

כפי שצוין לעיל, ברגע שמפתח נכתב למשבצת, הוא לא מוזז יותר. גם כאשר אלמנט נמחק מהטבלה, המפתח נשאר במקומו, הערך שלו פשוט הופך ריק. זה אומר שאנחנו לא צריכים להשתמש בפעולת כתיבה אטומית לערך החריץ, כי זה לא משנה אם הערך הנוכחי ריק או לא - הוא עדיין יהפוך לריק.

שינוי גודל טבלת גיבוב

ניתן לשנות את הגודל של טבלת hash על ידי יצירת טבלה גדולה יותר והכנסת אלמנטים לא ריקים מהטבלה הישנה לתוכה. לא יישמתי את הפונקציונליות הזו כי רציתי לשמור על קוד לדוגמה פשוט. יתרה מכך, בתוכניות CUDA, הקצאת זיכרון נעשית לרוב בקוד המארח ולא בליבת CUDA.

המאמר טבלת Hash ללא נעילה ללא המתנה מתאר כיצד לשנות מבנה נתונים כזה מוגן נעילה.

תחרותיות

בקטעי הקוד של הפונקציה לעיל 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);
    }
}

טבלת הגיבוב עמידה לנעילה תומכת בהוספות, חיפושים ומחיקות במקביל. מכיוון שצמדי מפתח-ערך נמצאים תמיד באחד מארבעה מצבים והמפתחות אינם זזים, הטבלה מבטיחה נכונות גם כאשר נעשה שימוש בסוגים שונים של פעולות בו-זמנית.

עם זאת, אם נעבד אצווה של הוספות ומחיקות במקביל, ואם מערך הקלט של הזוגות מכיל מפתחות כפולים, אז לא נוכל לחזות אילו זוגות "יזכו" - ייכתבו אחרון לטבלת ה-hash. נניח שקראנו לקוד ההכנסה עם מערך קלט של זוגות 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() השתמש במצביע רגיל או נדיף למערך של זוגות בטבלת ה-hash. תיעוד CUDA מציין ש:

המהדר עשוי לבחור לבצע אופטימיזציה של קריאה וכתיבה לזיכרון גלובלי או משותף... ניתן להשבית אופטימיזציות אלו באמצעות מילת המפתח volatile: ... כל התייחסות למשתנה זה מורכבת להוראת קריאה או כתיבה בזיכרון אמיתי.

שיקולי תקינות אינם מחייבים יישום volatile. אם שרשור הביצוע משתמש בערך שמור מפעולת קריאה קודמת, אז הוא ישתמש במידע מעט מיושן. אבל עדיין, זהו מידע מהמצב הנכון של טבלת ה-hash ברגע מסוים של קריאת הקרנל. אם אתה צריך להשתמש במידע העדכני ביותר, אתה יכול להשתמש באינדקס volatile, אבל אז הביצועים ירדו מעט: לפי הבדיקות שלי, בעת מחיקת 32 מיליון אלמנטים, המהירות ירדה מ-500 מיליון מחיקות/שנייה ל-450 מיליון מחיקות/שנייה.

פרודוקטיביות

במבחן להכנסת 64 מיליון אלמנטים ומחיקת 32 מיליון מהם, תחרות בין std::unordered_map ואין כמעט טבלת גיבוב עבור ה-GPU:

std::unordered_map בילה 70 אלפיות השנייה בהכנסה והסרה של אלמנטים ולאחר מכן בשחרורם unordered_map (להיפטר ממיליוני אלמנטים לוקח הרבה זמן, כי בפנים unordered_map מתבצעות הקצאות זיכרון מרובות). בכנות, std:unordered_map הגבלות שונות לחלוטין. זהו חוט הפעלה יחיד של מעבד, תומך בערכי מפתח בכל גודל, מתפקד היטב בשיעורי ניצול גבוהים ומציג ביצועים יציבים לאחר מחיקות מרובות.

משך טבלת הגיבוב עבור ה-GPU והתקשורת בין תוכניות הייתה 984 אלפיות השנייה. זה כולל את הזמן המושקע בהצבת הטבלה בזיכרון ומחיקתה (הקצאת 1 ג'יגה-בייט של זיכרון פעם אחת, מה שלוקח זמן מה ב-CUDA), הכנסת ומחיקת אלמנטים ואיטרציה עליהם. כל העותקים מזיכרון כרטיס המסך וממנו נלקחים בחשבון גם כן.

לטבלת הגיבוב עצמה נדרשו 271 אלפיות השנייה כדי להשלים. זה כולל את הזמן המושקע על ידי כרטיס המסך בהכנסת ומחיקת אלמנטים, ואינו לוקח בחשבון את הזמן המושקע בהעתקה לזיכרון ואיטרציה על הטבלה המתקבלת. אם טבלת ה-GPU קיימת לאורך זמן, או אם טבלת ה-hash כלולה בזיכרון של כרטיס המסך (לדוגמה, כדי ליצור טבלת גיבוב שתשמש את קוד ה-GPU אחר ולא את המעבד המרכזי), אז תוצאת הבדיקה רלוונטית.

טבלת הגיבוב לכרטיס מסך מדגימה ביצועים גבוהים הודות לתפוקה גבוהה והקבלה אקטיבית.

מגבלות

לארכיטקטורת טבלת הגיבוב יש כמה בעיות שכדאי לשים לב אליהן:

• חיטוט ליניארי נפגע על ידי clustering, מה שגורם למפתחות בטבלה להיות פחות ממושלם.
• מפתחות אינם מוסרים באמצעות הפונקציה delete ועם הזמן הם עומסים את השולחן.

כתוצאה מכך, הביצועים של טבלת גיבוב עלולים להתדרדר בהדרגה, במיוחד אם היא קיימת במשך זמן רב ויש לה הוספות ומחיקות רבות. אחת הדרכים לצמצם את החסרונות הללו היא להכניס מחדש לטבלה חדשה עם שיעור ניצול נמוך למדי ולסנן את המפתחות שהוסרו במהלך הגיבוש מחדש.

כדי להמחיש את הבעיות המתוארות, אשתמש בקוד שלמעלה כדי ליצור טבלה עם 128 מיליון אלמנטים ולעבור בלולאה של 4 מיליון אלמנטים עד שאמלא 124 מיליון משבצות (שיעור ניצול של בערך 0,96). הנה טבלת התוצאות, כל שורה היא קריאת ליבת CUDA להכנסת 4 מיליון אלמנטים חדשים לטבלת hash אחת:

שיעור שימוש
משך הכנסה 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 מיליון מפתחות/שנייה)

ככל שהניצול עולה, הביצועים יורדים. זה לא רצוי ברוב המקרים. אם יישום מכניס אלמנטים לטבלה ואז משליך אותם (לדוגמה, בעת ספירת מילים בספר), זו לא בעיה. אבל אם האפליקציה משתמשת בטבלת hash ארוכת חיים (לדוגמה, בעורך גרפי כדי לאחסן חלקים לא ריקים של תמונות שבהם המשתמש מוסיף ומוחק מידע לעתים קרובות), אז התנהגות זו עלולה להיות בעייתית.

ומדד את עומק החיטוט של טבלת ה-hash לאחר 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 ג'יגה-בייט), זו עדיין לא תהיה ההחלטה החכמה להפסיד כמויות כאלה.

בהמשך אכתוב עוד על יצירת טבלאות גיבוב לכרטיסי מסך שאין להם בעיות עם עומק החיטוט, כמו גם דרכים לשימוש חוזר בחריצים שנמחקו.

מדידת עומק צלילים

כדי לקבוע את עומק הבדיקה של מפתח, נוכל לחלץ את ה-hash של המפתח (אינדקס הטבלה האידיאלי שלו) מאינדקס הטבלה האמיתי שלו:

// 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.

מסקנה

אם יש לך שאלות או הערות, אנא שלח לי מייל לכתובת טויטר או לפתוח נושא חדש ב מאגרים.

קוד זה נכתב בהשראת מאמרים מצוינים:

• טבלת Hash הפשוטה בעולם ללא נעילה
• טבלת Hash ללא נעילה ללא המתנה

בעתיד, אמשיך לכתוב על יישומי טבלת hash עבור כרטיסי מסך ואנתח את הביצועים שלהם. התוכניות שלי כוללות שרשור, גיבוב רובין הוד וגיבוב קוקייה באמצעות פעולות אטומיות במבני נתונים ידידותיים ל-GPU.

מקור: www.habr.com