Tôi đã đăng nó trên Github .
Nó là một bảng băm GPU đơn giản có khả năng xử lý hàng trăm triệu lần chèn mỗi giây. Trên máy tính xách tay NVIDIA GTX 1060 của tôi, mã sẽ chèn 64 triệu cặp khóa-giá trị được tạo ngẫu nhiên trong khoảng 210 mili giây và loại bỏ 32 triệu cặp trong khoảng 64 mili giây.
Tức là tốc độ trên máy tính xách tay là khoảng 300 triệu lần chèn/giây và 500 triệu lần xóa/giây.
Bảng được viết bằng CUDA, mặc dù kỹ thuật tương tự có thể được áp dụng cho HLSL hoặc GLSL. Việc triển khai có một số hạn chế để đảm bảo hiệu suất cao trên card màn hình:
- Chỉ các khóa 32 bit và các giá trị giống nhau mới được xử lý.
- Bảng băm có kích thước cố định.
- Và kích thước này phải bằng hai lũy thừa.
Đối với khóa và giá trị, bạn cần đặt trước một dấu phân cách đơn giản (trong mã ở trên là 0xffffffff).
Bảng băm không có khóa
Bảng băm sử dụng địa chỉ mở với , nghĩa là, nó chỉ đơn giản là một mảng các cặp khóa-giá trị được lưu trữ trong bộ nhớ và có hiệu suất bộ đệm vượt trội. Điều tương tự không thể xảy ra đối với chuỗi, liên quan đến việc tìm kiếm một con trỏ trong danh sách liên kết. Bảng băm là một mảng lưu trữ các phần tử đơn giản KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
Kích thước của bảng là lũy thừa của hai chứ không phải số nguyên tố, vì một lệnh nhanh là đủ để áp dụng mặt nạ pow2/AND, nhưng toán tử mô đun chậm hơn nhiều. Điều này rất quan trọng trong trường hợp thăm dò tuyến tính, vì trong tra cứu bảng tuyến tính, chỉ mục vị trí phải được bao bọc trong mỗi vị trí. Và kết quả là chi phí của hoạt động được cộng thêm modulo vào mỗi vị trí.
Bảng chỉ lưu trữ khóa và giá trị cho từng phần tử chứ không lưu trữ hàm băm của khóa. Vì bảng chỉ lưu trữ các khóa 32 bit nên hàm băm được tính rất nhanh. Đoạn mã trên sử dụng hàm băm Murmur3, chỉ thực hiện một vài ca, XOR và phép nhân.
Bảng băm sử dụng các kỹ thuật khóa bảo vệ độc lập với thứ tự bộ nhớ. Ngay cả khi một số thao tác ghi làm gián đoạn thứ tự của các thao tác khác, bảng băm vẫn sẽ duy trì trạng thái chính xác. Chúng ta sẽ nói về điều này dưới đây. Kỹ thuật này hoạt động hiệu quả với các card màn hình chạy đồng thời hàng nghìn luồng.
Các khóa và giá trị trong bảng băm được khởi tạo để trống.
Mã có thể được sửa đổi để xử lý các khóa và giá trị 64 bit. Các khóa yêu cầu các hoạt động đọc, ghi và so sánh và hoán đổi nguyên tử. Và các giá trị yêu cầu thao tác đọc và ghi nguyên tử. May mắn thay, trong CUDA, các thao tác đọc-ghi đối với các giá trị 32 và 64 bit là nguyên tử miễn là chúng được căn chỉnh tự nhiên (xem bên dưới). ) và card màn hình hiện đại hỗ trợ các hoạt động so sánh và trao đổi nguyên tử 64-bit. Tất nhiên, khi chuyển sang 64 bit, hiệu năng sẽ giảm đi đôi chút.
Trạng thái bảng băm
Mỗi cặp khóa-giá trị trong bảng băm có thể có một trong bốn trạng thái:
- Khóa và giá trị trống. Ở trạng thái này, bảng băm được khởi tạo.
- Chìa khóa đã được ghi lại nhưng giá trị vẫn chưa được ghi lại. Nếu một luồng khác hiện đang đọc dữ liệu thì nó sẽ trả về trống. Điều này là bình thường, điều tương tự sẽ xảy ra nếu một luồng thực thi khác hoạt động sớm hơn một chút và chúng ta đang nói về cấu trúc dữ liệu đồng thời.
- Cả khóa và giá trị đều được ghi lại.
- Giá trị có sẵn cho các luồng thực thi khác, nhưng khóa thì chưa. Điều này có thể xảy ra do mô hình lập trình CUDA có mô hình bộ nhớ được sắp xếp lỏng lẻo. Điều này là bình thường; trong mọi trường hợp, khóa vẫn trống, ngay cả khi giá trị không còn trống nữa.
Một sắc thái quan trọng là khi khóa đã được ghi vào khe, nó sẽ không di chuyển nữa - ngay cả khi khóa bị xóa, chúng ta sẽ nói về điều này bên dưới.
Mã bảng băm thậm chí còn hoạt động với các mô hình bộ nhớ được sắp xếp lỏng lẻo trong đó không xác định được thứ tự đọc và ghi bộ nhớ. Khi chúng ta xem xét việc chèn, tra cứu và xóa trong bảng băm, hãy nhớ rằng mỗi cặp khóa-giá trị nằm ở một trong bốn trạng thái được mô tả ở trên.
Chèn vào bảng băm
Hàm CUDA chèn các cặp khóa-giá trị vào bảng băm trông như thế này:
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);
}
}
Để chèn một khóa, mã sẽ lặp qua mảng bảng băm bắt đầu bằng hàm băm của khóa được chèn. Mỗi vị trí trong mảng thực hiện thao tác so sánh và hoán đổi nguyên tử để so sánh khóa trong vị trí đó với khóa trống. Nếu phát hiện thấy không khớp, khóa trong vị trí sẽ được cập nhật bằng khóa được chèn và sau đó khóa vị trí ban đầu sẽ được trả về. Nếu khóa gốc này trống hoặc khớp với khóa được chèn thì mã sẽ tìm thấy một vị trí thích hợp để chèn và chèn giá trị được chèn vào vị trí đó.
Nếu trong một cuộc gọi kernel gpu_hashtable_insert() có nhiều phần tử có cùng khóa thì bất kỳ giá trị nào của chúng đều có thể được ghi vào khe khóa. Điều này được coi là bình thường: một trong các thao tác ghi khóa-giá trị trong cuộc gọi sẽ thành công, nhưng vì tất cả điều này xảy ra song song trong một số luồng thực thi nên chúng tôi không thể dự đoán việc ghi bộ nhớ nào sẽ là lần ghi cuối cùng.
Tra cứu bảng băm
Mã tìm kiếm khóa:
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);
}
}
Để tìm giá trị của khóa được lưu trữ trong bảng, chúng ta lặp qua mảng bắt đầu bằng hàm băm của khóa mà chúng ta đang tìm kiếm. Trong mỗi vị trí, chúng tôi kiểm tra xem khóa có phải là khóa chúng tôi đang tìm kiếm hay không và nếu có thì chúng tôi trả về giá trị của nó. Chúng tôi cũng kiểm tra xem khóa có trống không và nếu có thì chúng tôi sẽ hủy tìm kiếm.
Nếu chúng ta không thể tìm thấy khóa, mã sẽ trả về giá trị trống.
Tất cả các thao tác tìm kiếm này có thể được thực hiện đồng thời thông qua việc chèn và xóa. Mỗi cặp trong bảng sẽ có một trong bốn trạng thái được mô tả ở trên cho luồng.
Xóa trong bảng băm
Mã để xóa khóa:
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);
}
}
Việc xóa khóa được thực hiện theo cách khác thường: chúng ta để khóa trong bảng và đánh dấu giá trị của nó (không phải chính khóa) là trống. Mã này rất giống với lookup(), ngoại trừ việc khi tìm thấy một kết quả khớp trên một khóa, nó sẽ làm cho giá trị của nó trống.
Như đã đề cập ở trên, khi một khóa được ghi vào một khe, nó sẽ không còn được di chuyển nữa. Ngay cả khi một phần tử bị xóa khỏi bảng, khóa vẫn giữ nguyên, giá trị của nó sẽ trống. Điều này có nghĩa là chúng ta không cần sử dụng thao tác ghi nguyên tử cho giá trị vị trí, vì việc giá trị hiện tại có trống hay không không quan trọng - nó vẫn sẽ trống.
Thay đổi kích thước bảng băm
Bạn có thể thay đổi kích thước của bảng băm bằng cách tạo một bảng lớn hơn và chèn các phần tử không trống từ bảng cũ vào đó. Tôi không triển khai chức năng này vì tôi muốn giữ mã mẫu đơn giản. Hơn nữa, trong các chương trình CUDA, việc cấp phát bộ nhớ thường được thực hiện trong mã máy chủ hơn là trong nhân CUDA.
Trong bài viết mô tả cách sửa đổi cấu trúc dữ liệu được bảo vệ bằng khóa như vậy.
Năng lực cạnh tranh
Trong đoạn mã chức năng trên gpu_hashtable_insert(), _lookup() и _delete() xử lý một cặp khóa-giá trị tại một thời điểm. Và thấp hơn gpu_hashtable_insert(), _lookup() и _delete() xử lý song song một loạt các cặp, mỗi cặp trong một luồng thực thi GPU riêng biệt:
// 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);
}
}
Bảng băm chống khóa hỗ trợ chèn, tra cứu và xóa đồng thời. Vì các cặp khóa-giá trị luôn ở một trong bốn trạng thái và các khóa không di chuyển nên bảng đảm bảo tính chính xác ngay cả khi sử dụng đồng thời các loại thao tác khác nhau.
Tuy nhiên, nếu chúng ta xử lý song song một loạt các thao tác chèn và xóa và nếu mảng đầu vào của các cặp chứa các khóa trùng lặp thì chúng ta sẽ không thể dự đoán cặp nào sẽ “thắng”—sẽ được ghi vào bảng băm cuối cùng. Giả sử chúng ta gọi mã chèn với một mảng đầu vào gồm các cặp A/0 B/1 A/2 C/3 A/4. Khi mã hoàn thành, các cặp B/1 и C/3 được đảm bảo có mặt trong bảng, nhưng đồng thời bất kỳ cặp nào cũng sẽ xuất hiện trong đó A/0, A/2 hoặc A/4. Điều này có thể có hoặc không có vấn đề - tất cả phụ thuộc vào ứng dụng. Bạn có thể biết trước rằng không có khóa trùng lặp trong mảng đầu vào hoặc bạn có thể không quan tâm giá trị nào được ghi cuối cùng.
Nếu đây là vấn đề đối với bạn thì bạn cần tách các cặp trùng lặp thành các lệnh gọi hệ thống CUDA khác nhau. Trong CUDA, mọi thao tác gọi hạt nhân luôn hoàn thành trước lệnh gọi hạt nhân tiếp theo (ít nhất là trong một luồng. Trong các luồng khác nhau, hạt nhân được thực thi song song). Trong ví dụ trên, nếu bạn gọi một kernel bằng A/0 B/1 A/2 C/3, và cái kia với A/4, thì chìa khóa A sẽ nhận được giá trị 4.
Bây giờ hãy nói về việc liệu các chức năng có nên lookup() и delete() sử dụng một con trỏ đơn giản hoặc dễ bay hơi tới một mảng các cặp trong bảng băm. Tình trạng:
Trình biên dịch có thể chọn tối ưu hóa việc đọc và ghi vào bộ nhớ chung hoặc bộ nhớ dùng chung... Những tối ưu hóa này có thể bị tắt bằng cách sử dụng từ khóa
volatile: ... mọi tham chiếu đến biến này sẽ được biên dịch thành lệnh đọc hoặc ghi bộ nhớ thực.
Những cân nhắc về tính đúng đắn không cần phải áp dụng volatile. Nếu luồng thực thi sử dụng giá trị được lưu trong bộ nhớ đệm từ thao tác đọc trước đó thì luồng đó sẽ sử dụng thông tin hơi lỗi thời. Tuy nhiên, đây vẫn là thông tin từ trạng thái chính xác của bảng băm tại một thời điểm nhất định của lệnh gọi kernel. Nếu bạn cần sử dụng thông tin mới nhất, bạn có thể sử dụng chỉ mục volatile, nhưng sau đó hiệu suất sẽ giảm nhẹ: theo thử nghiệm của tôi, khi xóa 32 triệu phần tử, tốc độ giảm từ 500 triệu lần xóa/giây xuống còn 450 triệu lần xóa/giây.
Năng suất
Trong thử nghiệm chèn 64 triệu phần tử và xóa 32 triệu phần tử, sự cạnh tranh giữa std::unordered_map và hầu như không có bảng băm cho GPU:
std::unordered_map đã dành 70 ms để chèn và xóa các phần tử rồi giải phóng chúng unordered_map (việc loại bỏ hàng triệu phần tử mất rất nhiều thời gian, vì bên trong unordered_map nhiều phân bổ bộ nhớ được thực hiện). Nói một cách trung thực, std:unordered_map những hạn chế hoàn toàn khác nhau. Nó là một luồng thực thi CPU duy nhất, hỗ trợ các khóa-giá trị ở mọi kích thước, hoạt động tốt ở tốc độ sử dụng cao và cho thấy hiệu suất ổn định sau nhiều lần xóa.
Thời lượng của bảng băm cho GPU và giao tiếp giữa các chương trình là 984 mili giây. Điều này bao gồm thời gian dành cho việc đặt bảng vào bộ nhớ và xóa nó (phân bổ 1 GB bộ nhớ một lần, việc này mất một chút thời gian trong CUDA), chèn và xóa các phần tử cũng như lặp lại chúng. Tất cả các bản sao đến và đi từ bộ nhớ card màn hình cũng được tính đến.
Bản thân bảng băm mất 271 ms để hoàn thành. Điều này bao gồm thời gian dành cho thẻ video để chèn và xóa các phần tử, đồng thời không tính đến thời gian dành cho việc sao chép vào bộ nhớ và lặp lại bảng kết quả. Nếu bảng GPU tồn tại trong một thời gian dài hoặc nếu bảng băm được chứa hoàn toàn trong bộ nhớ của card màn hình (ví dụ: để tạo bảng băm sẽ được sử dụng bởi mã GPU khác chứ không phải bộ xử lý trung tâm), thì kết quả kiểm tra có liên quan.
Bảng băm cho card màn hình thể hiện hiệu suất cao do thông lượng cao và hoạt động song song tích cực.
Hạn chế
Kiến trúc bảng băm có một số vấn đề cần lưu ý:
- Việc thăm dò tuyến tính bị cản trở bởi việc phân cụm, khiến cho các phím trong bảng được đặt không được hoàn hảo.
- Các phím không bị xóa bằng chức năng
deletevà theo thời gian chúng làm bừa bộn bàn ăn.
Kết quả là hiệu suất của bảng băm có thể giảm dần, đặc biệt nếu nó tồn tại trong thời gian dài và có nhiều lần chèn và xóa. Một cách để giảm thiểu những nhược điểm này là thử lại vào một bảng mới với tỷ lệ sử dụng khá thấp và lọc ra các khóa đã bị loại bỏ trong quá trình thử lại.
Để minh họa các vấn đề được mô tả, tôi sẽ sử dụng đoạn mã trên để tạo một bảng có 128 triệu phần tử và lặp qua 4 triệu phần tử cho đến khi tôi lấp đầy 124 triệu vị trí (tỷ lệ sử dụng khoảng 0,96). Đây là bảng kết quả, mỗi hàng là một lệnh gọi kernel CUDA để chèn 4 triệu phần tử mới vào một bảng băm:
Tỷ lệ sử dụng
Thời lượng chèn 4 phần tử
0,00
11,608448 ms (361,314798 triệu phím/giây)
0,03
11,751424 ms (356,918799 triệu phím/giây)
0,06
11,942592 ms (351,205515 triệu phím/giây)
0,09
12,081120 ms (347,178429 triệu phím/giây)
0,12
12,242560 ms (342,600233 triệu phím/giây)
0,16
12,396448 ms (338,347235 triệu phím/giây)
0,19
12,533024 ms (334,660176 triệu phím/giây)
0,22
12,703328 ms (330,173626 triệu phím/giây)
0,25
12,884512 ms (325,530693 triệu phím/giây)
0,28
13,033472 ms (321,810182 triệu phím/giây)
0,31
13,239296 ms (316,807174 triệu phím/giây)
0,34
13,392448 ms (313,184256 triệu phím/giây)
0,37
13,624000 ms (307,861434 triệu phím/giây)
0,41
13,875520 ms (302,280855 triệu phím/giây)
0,44
14,126528 ms (296,909756 triệu phím/giây)
0,47
14,399328 ms (291,284699 triệu phím/giây)
0,50
14,690304 ms (285,515123 triệu phím/giây)
0,53
15,039136 ms (278,892623 triệu phím/giây)
0,56
15,478656 ms (270,973402 triệu phím/giây)
0,59
15,985664 ms (262,379092 triệu phím/giây)
0,62
16,668673 ms (251,627968 triệu phím/giây)
0,66
17,587200 ms (238,486174 triệu phím/giây)
0,69
18,690048 ms (224,413765 triệu phím/giây)
0,72
20,278816 ms (206,831789 triệu phím/giây)
0,75
22,545408 ms (186,038058 triệu phím/giây)
0,78
26,053312 ms (160,989275 triệu phím/giây)
0,81
31,895008 ms (131,503463 triệu phím/giây)
0,84
42,103294 ms (99,619378 triệu phím/giây)
0,87
61,849056 ms (67,815164 triệu phím/giây)
0,90
105,695999 ms (39,682713 triệu phím/giây)
0,94
240,204636 ms (17,461378 triệu phím/giây)
Khi mức sử dụng tăng lên, hiệu suất sẽ giảm. Điều này là không mong muốn trong hầu hết các trường hợp. Nếu một ứng dụng chèn các phần tử vào một bảng rồi loại bỏ chúng (ví dụ: khi đếm các từ trong một cuốn sách) thì đây không phải là vấn đề. Nhưng nếu ứng dụng sử dụng bảng băm tồn tại lâu dài (ví dụ: trong trình chỉnh sửa đồ họa để lưu trữ các phần hình ảnh không trống mà người dùng thường xuyên chèn và xóa thông tin), thì hành vi này có thể có vấn đề.
Và đo độ sâu thăm dò bảng băm sau 64 triệu lần chèn (hệ số sử dụng 0,5). Độ sâu trung bình là 0,4774, vì vậy hầu hết các phím đều ở vị trí tốt nhất có thể hoặc cách vị trí tốt nhất một vị trí. Độ sâu âm thanh tối đa là 60.
Sau đó, tôi đo độ sâu thăm dò trên một chiếc bàn có 124 triệu hạt dao (hệ số sử dụng 0,97). Độ sâu trung bình đã là 10,1757 và mức tối đa - 6474 (!!). Hiệu suất cảm biến tuyến tính giảm đáng kể ở mức sử dụng cao.
Tốt nhất là giữ tỷ lệ sử dụng bảng băm này ở mức thấp. Nhưng sau đó chúng ta tăng hiệu suất nhưng phải trả giá bằng việc tiêu thụ bộ nhớ. May mắn thay, trong trường hợp khóa và giá trị 32 bit, điều này có thể hợp lý. Nếu trong ví dụ trên, trong một bảng có 128 triệu phần tử, chúng tôi giữ hệ số sử dụng là 0,25 thì chúng tôi có thể đặt không quá 32 triệu phần tử trong đó và 96 triệu vị trí còn lại sẽ bị mất - 8 byte cho mỗi cặp , 768 MB bộ nhớ bị mất.
Xin lưu ý rằng chúng ta đang nói về việc mất bộ nhớ card màn hình, đây là tài nguyên quý giá hơn bộ nhớ hệ thống. Mặc dù hầu hết các card đồ họa máy tính để bàn hiện đại hỗ trợ CUDA đều có bộ nhớ ít nhất 4 GB (tại thời điểm viết bài, NVIDIA 2080 Ti có 11 GB), nhưng việc mất số lượng như vậy vẫn không phải là quyết định khôn ngoan nhất.
Sau này tôi sẽ viết thêm về cách tạo bảng băm cho các card màn hình không gặp vấn đề về độ sâu thăm dò, cũng như cách sử dụng lại các vị trí đã xóa.
Đo độ sâu âm thanh
Để xác định độ sâu thăm dò của một khóa, chúng ta có thể trích xuất hàm băm của khóa (chỉ mục bảng lý tưởng của nó) từ chỉ mục bảng thực tế của nó:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Do sự kỳ diệu của số nhị phân bù hai và thực tế là dung lượng của bảng băm là hai lũy thừa của hai, phương pháp này sẽ hoạt động ngay cả khi chỉ mục chính được chuyển đến đầu bảng. Hãy lấy một khóa được băm thành 1 nhưng được chèn vào vị trí 3. Sau đó, đối với một bảng có dung lượng 4, chúng ta nhận được (3 — 1) & 3, tương đương với 2.
Kết luận
Nếu bạn có thắc mắc hoặc ý kiến, xin vui lòng gửi email cho tôi tại hoặc mở một chủ đề mới trong.
Mã này được viết dưới cảm hứng từ các bài viết xuất sắc:
Trong tương lai, tôi sẽ tiếp tục viết về cách triển khai bảng băm cho thẻ video và phân tích hiệu suất của chúng. Kế hoạch của tôi bao gồm xâu chuỗi, băm Robin Hood và băm chim cu bằng cách sử dụng các phép toán nguyên tử trong cấu trúc dữ liệu thân thiện với GPU.
Nguồn: www.habr.com