من آن را در Github پست کردم پروژه جدید A Simple Hash Table GPU.

این یک جدول هش GPU ساده است که قادر به پردازش صدها میلیون درج در ثانیه است. در لپ تاپ NVIDIA GTX 1060 من، کد 64 میلیون جفت کلید-مقدار تصادفی تولید شده را در حدود 210 میلی ثانیه درج می کند و 32 میلیون جفت را در حدود 64 میلی ثانیه حذف می کند.

یعنی سرعت یک لپ تاپ تقریباً 300 میلیون درج در ثانیه و 500 میلیون حذف در ثانیه است.

جدول به زبان CUDA نوشته شده است، اگرچه همین تکنیک را می توان برای HLSL یا GLSL اعمال کرد. پیاده سازی چندین محدودیت برای اطمینان از عملکرد بالا در کارت گرافیک دارد:

فقط کلیدهای 32 بیتی و مقادیر مشابه پردازش می شوند.
جدول هش اندازه ثابتی دارد.
و این اندازه باید برابر با دو به توان باشد.

برای کلیدها و مقادیر، باید یک نشانگر جداکننده ساده رزرو کنید (در کد بالا این 0xffffffff است).

میز هش بدون قفل

جدول هش از آدرس دهی باز با استفاده می کند کاوش خطی، یعنی به سادگی آرایه ای از جفت های کلید-مقدار است که در حافظه ذخیره می شود و عملکرد کش برتری دارد. همین امر را نمی توان در مورد زنجیره سازی، که شامل جستجوی یک اشاره گر در یک لیست پیوندی است، گفت. جدول هش یک آرایه ساده است که عناصر را ذخیره می کند KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

اندازه جدول یک توان دو است نه یک عدد اول، زیرا یک دستور سریع برای اعمال ماسک pow2/AND کافی است، اما عملگر مدول بسیار کندتر است. این در مورد کاوش خطی مهم است، زیرا در جستجوی جدول خطی، شاخص شکاف باید در هر شکاف پیچیده شود. و در نتیجه هزینه عملیات به هر اسلات مدول اضافه می شود.

جدول فقط کلید و مقدار را برای هر عنصر ذخیره می کند، نه یک هش از کلید. از آنجایی که جدول فقط کلیدهای 32 بیتی را ذخیره می کند، هش خیلی سریع محاسبه می شود. کد بالا از هش Murmur3 استفاده می کند که فقط چند شیفت، XOR و ضرب را انجام می دهد.

جدول هش از تکنیک های محافظت قفل استفاده می کند که مستقل از ترتیب حافظه هستند. حتی اگر برخی از عملیات نوشتن نظم سایر عملیات‌ها را مختل کنند، جدول هش همچنان وضعیت صحیح را حفظ می‌کند. در ادامه در این مورد صحبت خواهیم کرد. این تکنیک با کارت‌های ویدیویی که هزاران رشته را به طور همزمان اجرا می‌کنند عالی عمل می‌کند.

کلیدها و مقادیر در جدول هش مقداردهی اولیه می شوند تا خالی شوند.

کد را می توان برای کنترل کلیدها و مقادیر 64 بیتی نیز تغییر داد. کلیدها به عملیات خواندن، نوشتن، و مقایسه و تعویض اتمی نیاز دارند. و مقادیر به عملیات خواندن و نوشتن اتمی نیاز دارند. خوشبختانه، در CUDA، عملیات خواندن-نوشتن برای مقادیر 32 و 64 بیتی تا زمانی که به طور طبیعی تراز شده باشند، اتمی هستند (به زیر مراجعه کنید). اینجا) و کارت های ویدئویی مدرن از عملیات مقایسه و تبادل اتمی 64 بیتی پشتیبانی می کنند. البته هنگام انتقال به 64 بیت، عملکرد کمی کاهش می یابد.

وضعیت جدول هش

هر جفت کلید-مقدار در جدول هش می تواند یکی از چهار حالت را داشته باشد:

کلید و مقدار خالی هستند. در این حالت جدول هش مقداردهی اولیه می شود.
کلید نوشته شده است، اما مقدار آن هنوز یادداشت نشده است. اگر رشته دیگری در حال حاضر داده ها را می خواند، سپس خالی برمی گردد. این طبیعی است، اگر یک رشته اجرایی دیگر کمی زودتر کار می کرد، همین اتفاق می افتاد و ما در مورد ساختار داده همزمان صحبت می کنیم.
هم کلید و هم مقدار ثبت می شود.
مقدار برای رشته های دیگر اجرا در دسترس است، اما کلید هنوز موجود نیست. این می تواند به این دلیل اتفاق بیفتد که مدل برنامه نویسی CUDA دارای یک مدل حافظه منظم است. این طبیعی است؛ در هر صورت، کلید همچنان خالی است، حتی اگر مقدار دیگر خالی باشد.

نکته مهم این است که وقتی کلید در شکاف نوشته شد، دیگر حرکت نمی کند - حتی اگر کلید حذف شود، در زیر در مورد این موضوع صحبت خواهیم کرد.

کد جدول هش حتی با مدل‌های حافظه‌ای که مرتب شده‌اند کار می‌کند که ترتیب خواندن و نوشتن حافظه در آن‌ها مشخص نیست. همانطور که به درج، جستجو و حذف در جدول هش نگاه می کنیم، به یاد داشته باشید که هر جفت کلید-مقدار در یکی از چهار حالتی است که در بالا توضیح داده شد.

درج در جدول هش

تابع CUDA که جفت‌های کلید-مقدار را در جدول هش قرار می‌دهد به شکل زیر است:

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

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

برای درج یک کلید، کد از طریق آرایه جدول هش با هش کلید درج شده شروع می شود. هر شکاف در آرایه یک عملیات مقایسه و تعویض اتمی را انجام می دهد که کلید موجود در آن شکاف را با خالی مقایسه می کند. در صورت تشخیص عدم تطابق، کلید موجود در شکاف با کلید درج شده به روز می شود و سپس کلید شکاف اصلی برگردانده می شود. اگر این کلید اصلی خالی بود یا با کلید درج شده مطابقت داشت، کد یک شکاف مناسب برای درج پیدا کرد و مقدار درج شده را در شکاف وارد کرد.

اگر در یک هسته فراخوانی شود gpu_hashtable_insert() چندین عنصر با یک کلید وجود دارد، سپس هر یک از مقادیر آنها را می توان در شکاف کلید نوشت. این امر طبیعی در نظر گرفته می‌شود: یکی از نوشتن‌های کلید-مقدار در طول تماس با موفقیت انجام می‌شود، اما از آنجایی که همه اینها به صورت موازی در چندین رشته اجرا اتفاق می‌افتد، نمی‌توانیم پیش‌بینی کنیم که کدام حافظه آخرین نوشته باشد.

جستجوی جدول هش

کد برای جستجوی کلیدها:

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

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

برای یافتن مقدار یک کلید ذخیره شده در جدول، از طریق آرایه ای که با هش کلید مورد نظرمان شروع می شود، تکرار می کنیم. در هر شکاف بررسی می‌کنیم که آیا کلید همانی است که به دنبال آن هستیم یا نه، و اگر چنین است، مقدار آن را برمی‌گردانیم. همچنین بررسی می کنیم که آیا کلید خالی است یا خیر، و اگر چنین است، جستجو را لغو می کنیم.

اگر نتوانیم کلید را پیدا کنیم، کد یک مقدار خالی برمی گرداند.

همه این عملیات جستجو را می توان همزمان از طریق درج و حذف انجام داد. هر جفت در جدول دارای یکی از چهار حالتی است که در بالا برای جریان توضیح داده شد.

حذف در جدول هش

کد حذف کلیدها:

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

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

حذف یک کلید به روشی غیرعادی انجام می شود: کلید را در جدول می گذاریم و مقدار آن (نه خود کلید) را خالی علامت گذاری می کنیم. این کد بسیار شبیه به lookup()، با این تفاوت که وقتی مطابقی روی یک کلید پیدا می شود، مقدار آن را خالی می کند.

همانطور که در بالا ذکر شد، هنگامی که یک کلید در یک شکاف نوشته می شود، دیگر منتقل نمی شود. حتی زمانی که یک عنصر از جدول حذف می شود، کلید در جای خود باقی می ماند، مقدار آن به سادگی خالی می شود. این بدان معنی است که ما نیازی به استفاده از عملیات نوشتن اتمی برای مقدار شکاف نداریم، زیرا مهم نیست که مقدار فعلی خالی باشد یا نه - همچنان خالی خواهد بود.

تغییر اندازه جدول هش

می توانید اندازه جدول هش را با ایجاد یک جدول بزرگتر و قرار دادن عناصر غیر خالی از جدول قدیمی در آن تغییر دهید. من این قابلیت را پیاده سازی نکردم زیرا می خواستم کد نمونه را ساده نگه دارم. علاوه بر این، در برنامه های CUDA، تخصیص حافظه اغلب در کد میزبان انجام می شود تا در هسته CUDA.

مقاله جدول هش انتظار بدون قفل نحوه اصلاح چنین ساختار داده ای محافظت شده با قفل را شرح می دهد.

رقابت پذیری

در قطعه کد تابع بالا gpu_hashtable_insert(), _lookup() и _delete() پردازش یک جفت کلید-مقدار در یک زمان. و پایین تر gpu_hashtable_insert(), _lookup() и _delete() آرایه ای از جفت ها را به صورت موازی پردازش کنید، هر جفت در یک رشته اجرای GPU جداگانه:

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

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

جدول هش مقاوم در برابر قفل از درج‌ها، جستجوها و حذف‌های همزمان پشتیبانی می‌کند. از آنجایی که جفت‌های کلید-مقدار همیشه در یکی از چهار حالت هستند و کلیدها حرکت نمی‌کنند، جدول صحت را تضمین می‌کند حتی زمانی که انواع مختلف عملیات به طور همزمان استفاده می‌شوند.

با این حال، اگر دسته‌ای از درج‌ها و حذف‌ها را به صورت موازی پردازش کنیم، و اگر آرایه ورودی جفت‌ها حاوی کلیدهای تکراری باشد، نمی‌توانیم پیش‌بینی کنیم که کدام جفت‌ها «برنده» خواهند شد - آخرین بار در جدول هش نوشته می‌شود. فرض کنید کد درج را با یک آرایه ورودی از جفت فراخوانی کردیم A/0 B/1 A/2 C/3 A/4. وقتی کد کامل شد، جفت می شود B/1 и C/3 تضمین می شود که در جدول حضور داشته باشند، اما در همان زمان هر یک از جفت ها در آن ظاهر می شوند A/0, A/2 یا A/4. این ممکن است یک مشکل باشد یا نباشد - همه چیز به برنامه بستگی دارد. ممکن است از قبل بدانید که هیچ کلید تکراری در آرایه ورودی وجود ندارد، یا ممکن است اهمیتی نداشته باشید که آخرین مقدار کدام مقدار نوشته شده است.

اگر این مشکل برای شماست، باید جفت های تکراری را به فراخوانی های مختلف سیستم CUDA جدا کنید. در CUDA، هر عملیاتی که هسته را فراخوانی می کند، همیشه قبل از فراخوانی هسته بعدی کامل می شود (حداقل در یک رشته. در رشته های مختلف، هسته ها به صورت موازی اجرا می شوند). در مثال بالا، اگر یک کرنل را با آن فراخوانی کنید A/0 B/1 A/2 C/3، و دیگری با A/4، سپس کلید A ارزش را دریافت خواهد کرد 4.

حالا بیایید در مورد اینکه آیا توابع باید صحبت کنیم lookup() и delete() از یک اشاره گر ساده یا فرار به آرایه ای از جفت ها در جدول هش استفاده کنید. اسناد CUDA کشورهایی که:

کامپایلر ممکن است خواندن و نوشتن را در حافظه عمومی یا مشترک بهینه سازی کند... این بهینه سازی ها را می توان با استفاده از کلمه کلیدی غیرفعال کرد. volatile: ... هر ارجاعی به این متغیر در یک دستورالعمل خواندن یا نوشتن حافظه واقعی کامپایل می شود.

ملاحظات صحت نیازی به کاربرد ندارد volatile. اگر رشته اجرا از یک مقدار ذخیره شده از یک عملیات خواندن قبلی استفاده کند، از اطلاعات کمی قدیمی استفاده می کند. اما با این حال، این اطلاعات از وضعیت صحیح جدول هش در یک لحظه خاص از فراخوانی هسته است. اگر نیاز به استفاده از آخرین اطلاعات دارید، می توانید از فهرست استفاده کنید volatile، اما پس از آن عملکرد کمی کاهش می یابد: طبق آزمایشات من، هنگام حذف 32 میلیون عنصر، سرعت از 500 میلیون حذف در ثانیه به 450 میلیون حذف در ثانیه کاهش یافت.

کارایی

در آزمون درج 64 میلیون عنصر و حذف 32 میلیون از آنها، رقابت بین std::unordered_map و تقریبا هیچ جدول هش برای GPU وجود ندارد:

std::unordered_map 70 میلی‌ثانیه را صرف قرار دادن و حذف عناصر و سپس آزاد کردن آنها کرد unordered_map (خلاص شدن از شر میلیون ها عنصر زمان زیادی می برد، زیرا در داخل unordered_map تخصیص حافظه چندگانه انجام می شود). صادقانه بگویم، std:unordered_map محدودیت های کاملا متفاوت این یک رشته اجرایی واحد CPU است، از مقادیر کلیدی با هر اندازه پشتیبانی می‌کند، در نرخ‌های استفاده بالا عملکرد خوبی دارد و پس از حذف‌های متعدد، عملکرد پایداری را نشان می‌دهد.

مدت زمان جدول هش برای GPU و ارتباطات بین برنامه ای 984 میلی ثانیه بود. این شامل زمان صرف شده برای قرار دادن جدول در حافظه و حذف آن (تخصیص 1 گیگابایت حافظه برای یک بار، که مدتی در CUDA طول می کشد)، درج و حذف عناصر و تکرار روی آنها می باشد. تمام کپی ها به و از حافظه کارت ویدیو نیز در نظر گرفته می شود.

خود جدول هش 271 میلی ثانیه طول کشید تا تکمیل شود. این شامل زمان صرف شده توسط کارت ویدیو برای درج و حذف عناصر است و زمان صرف شده برای کپی کردن در حافظه و تکرار در جدول حاصل را در نظر نمی گیرد. اگر جدول GPU برای مدت طولانی زندگی می کند، یا اگر جدول هش به طور کامل در حافظه کارت گرافیک وجود دارد (به عنوان مثال، برای ایجاد یک جدول هش که توسط سایر کدهای GPU و نه پردازنده مرکزی استفاده می شود)، سپس نتیجه آزمایش مرتبط است

جدول هش برای کارت گرافیک به دلیل توان عملیاتی بالا و موازی سازی فعال، عملکرد بالایی را نشان می دهد.

محدودیت ها

معماری جدول هش چند مسئله دارد که باید از آنها آگاه بود:

کاوش خطی به دلیل خوشه بندی با مشکل مواجه می شود، که باعث می شود کلیدها در جدول کمتر از حد کامل قرار بگیرند.
کلیدها با استفاده از عملکرد حذف نمی شوند delete و به مرور زمان میز را به هم می ریزند.

در نتیجه، عملکرد یک جدول هش می تواند به تدریج کاهش یابد، به خصوص اگر برای مدت طولانی وجود داشته باشد و دارای درج ها و حذف های متعدد باشد. یکی از راه‌های کاهش این معایب این است که در جدول جدید با نرخ استفاده نسبتاً پایین، مجدداً هش کنید و کلیدهای حذف شده را در حین هش کردن مجدد فیلتر کنید.

برای نشان دادن مسائلی که توضیح داده شد، از کد بالا برای ایجاد جدولی با 128 میلیون عنصر و حلقه زدن 4 میلیون عنصر استفاده می کنم تا زمانی که 124 میلیون اسلات را پر کنم (نرخ استفاده حدود 0,96). در اینجا جدول نتیجه است، هر ردیف یک فراخوانی هسته CUDA برای درج 4 میلیون عنصر جدید در یک جدول هش است:

نرخ بهره برداری
مدت زمان درج 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 مگابایت حافظه از دست رفته.

لطفاً توجه داشته باشید که ما در مورد از بین رفتن حافظه کارت گرافیک صحبت می کنیم که منبع ارزشمندتری نسبت به حافظه سیستم است. اگرچه اکثر کارت‌های گرافیک مدرن رومیزی که از CUDA پشتیبانی می‌کنند حداقل 4 گیگابایت حافظه دارند (در زمان نگارش، NVIDIA 2080 Ti دارای 11 گیگابایت حافظه است)، باز هم عاقلانه‌ترین تصمیم برای از دست دادن چنین مقادیری نخواهد بود.

بعداً در مورد ایجاد جداول هش برای کارت‌های ویدیویی که مشکلی با عمق کاوش ندارند و همچنین روش‌های استفاده مجدد از اسلات‌های حذف شده بیشتر خواهم نوشت.

اندازه گیری عمق صدا

برای تعیین عمق کاوش یک کلید، می توانیم هش کلید (شاخص جدول ایده آل آن) را از شاخص جدول واقعی آن استخراج کنیم:

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

به دلیل جادوی اعداد باینری مکمل دو و این واقعیت که ظرفیت جدول هش دو به توان دو است، این رویکرد حتی زمانی که شاخص کلید به ابتدای جدول منتقل شود، کار خواهد کرد. بیایید کلیدی را بگیریم که به 1 هش شده است، اما در شکاف 3 قرار داده شده است. سپس برای جدولی با ظرفیت 4، دریافت می کنیم (3 — 1) & 3، که معادل 2 است.

نتیجه

اگر سؤال یا نظری دارید، لطفاً به من ایمیل بزنید توییتر یا یک موضوع جدید باز کنید مخازن.

این کد با الهام از مقالات عالی نوشته شده است:

در آینده، من به نوشتن در مورد پیاده سازی جدول هش برای کارت های ویدئویی و تجزیه و تحلیل عملکرد آنها ادامه خواهم داد. برنامه‌های من شامل زنجیره‌سازی، هش کردن رابین هود و هش کردن فاخته با استفاده از عملیات اتمی در ساختارهای داده‌ای است که سازگار با GPU هستند.

منبع: www.habr.com

جدول هش ساده برای GPU