Github에 올렸어요 새 프로젝트 간단한 GPU 해시 테이블.

초당 수억 개의 삽입을 처리할 수 있는 간단한 GPU 해시 테이블입니다. 내 NVIDIA GTX 1060 노트북에서 코드는 약 64ms 안에 무작위로 생성된 210만 개의 키-값 쌍을 삽입하고 약 32ms 안에 64만 쌍을 제거합니다.

즉, 랩탑의 속도는 초당 삽입 300억 회, 삭제 500억 회/초입니다.

테이블은 CUDA로 작성되었지만 HLSL이나 GLSL에도 동일한 기술을 적용할 수 있습니다. 구현에는 비디오 카드의 고성능을 보장하기 위한 몇 가지 제한 사항이 있습니다.

• 32비트 키와 동일한 값만 처리됩니다.
• 해시 테이블의 크기는 고정되어 있습니다.
• 그리고 이 크기는 XNUMX의 거듭제곱과 같아야 합니다.

키와 값의 경우 간단한 구분 기호를 예약해야 합니다(위 코드에서는 0xffffffff임).

잠금이 없는 해시 테이블

해시 테이블은 다음과 같은 개방형 주소 지정을 사용합니다. 선형 프로빙즉, 단순히 메모리에 저장되고 우수한 캐시 성능을 갖는 키-값 쌍의 배열입니다. 연결된 목록에서 포인터를 검색하는 연결의 경우에도 마찬가지입니다. 해시 테이블은 요소를 저장하는 간단한 배열입니다. KeyValue:

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

테이블의 크기는 소수가 아닌 2의 거듭제곱입니다. 왜냐하면 하나의 빠른 명령어로 powXNUMX/AND 마스크를 적용하는 데 충분하지만 모듈러스 연산자는 훨씬 느리기 때문입니다. 이는 선형 프로빙의 경우 중요합니다. 선형 테이블 조회에서는 슬롯 인덱스가 각 슬롯에 래핑되어야 하기 때문입니다. 결과적으로 작업 비용은 각 슬롯에 모듈로 추가됩니다.

테이블에는 키의 해시가 아닌 각 요소의 키와 값만 저장됩니다. 테이블에는 32비트 키만 저장되므로 해시가 매우 빠르게 계산됩니다. 위의 코드는 몇 번의 교대, XOR 및 곱셈만 수행하는 Murmur3 해시를 사용합니다.

해시 테이블은 메모리 순서와 무관한 잠금 보호 기술을 사용합니다. 일부 쓰기 작업이 다른 작업의 순서를 방해하더라도 해시 테이블은 여전히 ​​올바른 상태를 유지합니다. 이에 대해서는 아래에서 이야기하겠습니다. 이 기술은 수천 개의 스레드를 동시에 실행하는 비디오 카드에서 훌륭하게 작동합니다.

해시 테이블의 키와 값은 비어 있도록 초기화됩니다.

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억 XNUMX천만 삭제로 감소했습니다.

Производительность

64만개 요소를 삽입하고 32만개를 삭제하는 테스트에서 std::unordered_map GPU에는 사실상 해시 테이블이 없습니다.

std::unordered_map 요소를 삽입하고 제거한 후 해제하는 데 70ms가 소요되었습니다. unordered_map (수백만 개의 요소를 제거하는 데는 많은 시간이 걸립니다. unordered_map 다중 메모리 할당이 이루어집니다). 솔직히 말하자면, std:unordered_map 완전히 다른 제한 사항. 단일 CPU 실행 스레드로, 모든 크기의 키-값을 지원하고, 높은 활용률에서도 좋은 성능을 발휘하며, 여러 번 삭제 후에도 안정적인 성능을 보여줍니다.

GPU 및 프로그램 간 통신을 위한 해시 테이블의 지속 시간은 984ms였습니다. 여기에는 테이블을 메모리에 배치하고 삭제하는 데 소요되는 시간(한 번에 1GB의 메모리를 할당하는 데 CUDA에서는 약간의 시간이 소요됨), 요소를 삽입 및 삭제하고 반복하는 데 소요되는 시간이 포함됩니다. 비디오 카드 메모리로 들어오고 나가는 모든 복사본도 고려됩니다.

해시 테이블 자체를 완료하는 데 271ms가 걸렸습니다. 여기에는 비디오 카드에서 요소를 삽입하고 삭제하는 데 소요된 시간이 포함되며, 메모리에 복사하고 결과 테이블을 반복하는 데 소요된 시간은 고려되지 않습니다. GPU 테이블이 오랫동안 지속되거나 해시 테이블이 비디오 카드의 메모리에 완전히 포함되어 있는 경우(예: 중앙 프로세서가 아닌 다른 GPU 코드에서 사용할 해시 테이블을 생성하려는 경우) 테스트 결과가 관련이 있습니다.

비디오 카드의 해시 테이블은 높은 처리량과 활성 병렬화로 인해 높은 성능을 보여줍니다.

제한

해시 테이블 아키텍처에는 알아야 할 몇 가지 문제가 있습니다.

• 선형 프로빙은 클러스터링으로 인해 방해를 받으며, 이로 인해 테이블의 키가 완벽하지 않게 배치됩니다.
• 기능을 사용해도 키가 제거되지 않습니다. delete 시간이 지남에 따라 테이블이 어수선해집니다.

결과적으로 해시 테이블의 성능은 점차 저하될 수 있습니다. 특히 오랫동안 존재하고 삽입 및 삭제가 많은 경우에는 더욱 그렇습니다. 이러한 단점을 완화하는 한 가지 방법은 활용률이 상당히 낮은 새 테이블로 다시 해시하고 다시 해시 중에 제거된 키를 필터링하는 것입니다.

설명된 문제를 설명하기 위해 위 코드를 사용하여 128억 4천 124백만 개의 요소가 있는 테이블을 만들고 0,96억 4천 XNUMX백만 개의 슬롯을 채울 때까지 XNUMX백만 개의 요소를 반복합니다(사용률 약 XNUMX). 다음은 결과 테이블입니다. 각 행은 하나의 해시 테이블에 XNUMX만 개의 새 요소를 삽입하는 CUDA 커널 호출입니다.

이용률
삽입 기간 4개 요소

0,00
11,608448ms(361,314798억XNUMX만XNUMX만 키/초)

0,03
11,751424ms(356,918799억XNUMX만XNUMX만 키/초)

0,06
11,942592ms(351,205515억XNUMX만XNUMX만 키/초)

0,09
12,081120ms(347,178429억XNUMX만XNUMX만 키/초)

0,12
12,242560ms(342,600233억XNUMX만XNUMX만 키/초)

0,16
12,396448ms(338,347235억XNUMX만XNUMX만 키/초)

0,19
12,533024ms(334,660176억XNUMX만XNUMX만 키/초)

0,22
12,703328ms(330,173626억XNUMX만XNUMX만 키/초)

0,25
12,884512ms(325,530693억XNUMX만XNUMX만 키/초)

0,28
13,033472ms(321,810182억XNUMX만XNUMX만 키/초)

0,31
13,239296ms(316,807174억XNUMX만XNUMX만 키/초)

0,34
13,392448ms(313,184256억XNUMX만XNUMX만 키/초)

0,37
13,624000ms(307,861434억XNUMX만XNUMX만 키/초)

0,41
13,875520ms(302,280855억XNUMX만XNUMX만 키/초)

0,44
14,126528ms(296,909756억XNUMX만XNUMX만 키/초)

0,47
14,399328ms(291,284699억XNUMX만XNUMX만 키/초)

0,50
14,690304ms(285,515123억XNUMX만XNUMX만 키/초)

0,53
15,039136ms(278,892623억XNUMX만XNUMX만 키/초)

0,56
15,478656ms(270,973402억XNUMX만XNUMX만 키/초)

0,59
15,985664ms(262,379092억XNUMX만XNUMX만 키/초)

0,62
16,668673ms(251,627968억XNUMX만XNUMX만 키/초)

0,66
17,587200ms(238,486174억XNUMX만XNUMX만 키/초)

0,69
18,690048ms(224,413765억XNUMX만XNUMX만 키/초)

0,72
20,278816ms(206,831789억XNUMX만XNUMX만 키/초)

0,75
22,545408ms(186,038058억XNUMX만XNUMX만 키/초)

0,78
26,053312ms(160,989275억XNUMX만XNUMX만 키/초)

0,81
31,895008ms(131,503463억XNUMX만XNUMX만 키/초)

0,84
42,103294ms(99,619378억XNUMX만XNUMX만 키/초)

0,87
61,849056ms(67,815164억XNUMX만XNUMX만 키/초)

0,90
105,695999ms(39,682713억XNUMX만XNUMX만 키/초)

0,94
240,204636ms(17,461378억XNUMX만XNUMX만 키/초)

활용도가 증가하면 성능이 저하됩니다. 이는 대부분의 경우 바람직하지 않습니다. 응용 프로그램이 테이블에 요소를 삽입한 다음 해당 요소를 삭제하는 경우(예: 책의 단어 수를 셀 때) 이는 문제가 되지 않습니다. 그러나 응용 프로그램이 수명이 긴 해시 테이블을 사용하는 경우(예를 들어 사용자가 정보를 자주 삽입하고 삭제하는 이미지의 비어 있지 않은 부분을 저장하기 위해 그래픽 편집기에서) 이 동작은 문제가 될 수 있습니다.

그리고 64만 개의 삽입 후 해시 테이블 탐색 깊이를 측정했습니다(활용률 0,5). 평균 깊이는 0,4774이므로 대부분의 키는 가능한 가장 좋은 슬롯에 있거나 가장 좋은 위치에서 한 슬롯 떨어져 있습니다. 최대 소리 깊이는 60이었습니다.

그런 다음 124억 0,97만 개의 인서트(이용률 10,1757)가 있는 테이블에서 프로빙 깊이를 측정했습니다. 평균 깊이는 이미 XNUMX이었고 최대값은 - 6474 (!!). 선형 감지 성능은 활용도가 높을수록 크게 떨어집니다.

이 해시 테이블의 활용률을 낮게 유지하는 것이 가장 좋습니다. 그러나 메모리 소비를 희생하면서 성능을 향상시킵니다. 다행히 32비트 키와 값의 경우에는 이것이 정당화될 수 있습니다. 위의 예에서 128억 0,25만 개의 요소가 있는 테이블에서 활용률을 32로 유지하면 96만 개 이하의 요소를 배치할 수 있으며 나머지 8만 개의 슬롯은 손실됩니다(각 쌍당 768바이트). , XNUMXMB의 메모리 손실.

우리는 시스템 메모리보다 더 귀중한 리소스인 비디오 카드 메모리 손실에 대해 이야기하고 있습니다. CUDA를 지원하는 대부분의 최신 데스크톱 그래픽 카드에는 최소 4GB의 메모리가 있지만(작성 당시 NVIDIA 2080 Ti의 메모리는 11GB), 그러한 양을 잃는 것은 여전히 ​​가장 현명한 결정이 아닙니다.

나중에 프로빙 깊이에 문제가 없는 비디오 카드용 해시 테이블을 만드는 방법과 삭제된 슬롯을 재사용하는 방법에 대해 자세히 설명하겠습니다.

소리 깊이 측정

키의 검색 깊이를 결정하기 위해 실제 테이블 인덱스에서 키의 해시(이상적인 테이블 인덱스)를 추출할 수 있습니다.

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

1의 3의 보수 이진수의 마법과 해시 테이블의 용량이 4의 XNUMX제곱이라는 사실 때문에 이 접근 방식은 키 인덱스가 테이블의 시작 부분으로 이동하는 경우에도 작동합니다. XNUMX로 해시되었지만 슬롯 XNUMX에 삽입된 키를 가져옵니다. 그런 다음 용량이 XNUMX인 테이블에 대해 다음을 얻습니다. (3 — 1) & 3, 이는 2와 동일합니다.

결론

질문이나 의견이 있으시면 저에게 이메일을 보내주세요. 트위터 또는 다음에서 새 주제를 엽니다. 저장소.

이 코드는 훌륭한 기사에서 영감을 받아 작성되었습니다.

• 세계에서 가장 간단한 잠금 없는 해시 테이블
• 잠금이 없고 대기가 없는 해시 테이블

앞으로도 비디오 카드의 해시 테이블 구현에 대한 글을 쓰고 성능을 분석할 예정입니다. 내 계획에는 GPU 친화적인 데이터 구조에서 원자 연산을 사용한 체이닝, 로빈후드 해싱, 뻐꾸기 해싱이 포함됩니다.

출처 : habr.com