నేను దానిని 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 హాష్‌ని ఉపయోగిస్తుంది, ఇది కొన్ని షిఫ్ట్‌లు, 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 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