并行基数排序，这个实现实际上是如何工作的?有一些启发式方法吗? [英] Parallel radix sort, how would this implementation actually work? Are there some heuristics?
问题描述
我正在为他们的并行编程课程做一个 Udacity 测验.我很困惑我应该如何开始作业，因为我不确定我是否理解正确.
I am working on an Udacity quiz for their parallel programming course. I am pretty stuck on how I should start on the assignment because I am not sure if I understand it correctly.
对于赋值(在代码中)，我们得到了两个数组和值数组以及一个位置数组.我们应该使用并行基数排序对值数组进行排序，同时正确设置位置.
For the assignment (in code) we are given two arrays and array on values and an array of positions. We are supposed to sort the array of values with a parallelized radix sort, along with setting the positions correctly too.
我完全了解基数排序及其工作原理.我不明白的是他们希望我们如何实施它.这是开始分配的模板
I completely understand radix sort and how it works. What I don't understand is how they want us to implemented it. Here is the template given to start the assignment
//Udacity HW 4
//Radix Sorting
#include "reference_calc.cpp"
#include "utils.h"
/* Red Eye Removal
===============
For this assignment we are implementing red eye removal. This is
accomplished by first creating a score for every pixel that tells us how
likely it is to be a red eye pixel. We have already done this for you - you
are receiving the scores and need to sort them in ascending order so that we
know which pixels to alter to remove the red eye.
Note: ascending order == smallest to largest
Each score is associated with a position, when you sort the scores, you must
also move the positions accordingly.
Implementing Parallel Radix Sort with CUDA
==========================================
The basic idea is to construct a histogram on each pass of how many of each
"digit" there are. Then we scan this histogram so that we know where to put
the output of each digit. For example, the first 1 must come after all the
0s so we have to know how many 0s there are to be able to start moving 1s
into the correct position.
1) Histogram of the number of occurrences of each digit
2) Exclusive Prefix Sum of Histogram
3) Determine relative offset of each digit
For example [0 0 1 1 0 0 1]
-> [0 1 0 1 2 3 2]
4) Combine the results of steps 2 & 3 to determine the final
output location for each element and move it there
LSB Radix sort is an out-of-place sort and you will need to ping-pong values
between the input and output buffers we have provided. Make sure the final
sorted results end up in the output buffer! Hint: You may need to do a copy
at the end.
*/
void your_sort(unsigned int* const d_inputVals,
unsigned int* const d_inputPos,
unsigned int* const d_outputVals,
unsigned int* const d_outputPos,
const size_t numElems)
{
}
我特别不明白这 4 个步骤最终是如何对数组进行排序的.
I specifically don't understand how those 4 steps end up sorting the array.
所以第一步，我应该创建一个数字"的直方图(为什么用引号......?).因此，给定一个输入值 n，我需要将 0 和 1 计数成直方图.那么，第 1 步是否应该创建一个直方图数组，每个输入值一个?
So for the first step, I am supposed to create a histogram of the "digits" (why is that in quotes..?). So given a input value n I need to make a count of the 0's and 1's into a histogram. So, should step 1 create an array of histograms, one for each input value?
好吧，对于其余的步骤，它很快就会崩溃.有人可以告诉我这些步骤应该如何实现基数排序吗?
And well, for the rest of the steps it breaks down pretty quickly. Could someone show me how these steps are supposed to implement a radix sort?
推荐答案
基数排序背后的基本思想是，我们将考虑将每个元素逐位排序，从最不重要到最重要.对于每个数字，我们将移动元素，以使 这些数字按升序排列.
The basic idea behind a radix sort is that we will consider each element to be sorted digit by digit, from least significant to most significant. For each digit, we will move the elements so that those digits are in increasing order.
让我们举一个非常简单的例子.让我们对四个数量进行排序，每个数量都有 4 个二进制数字.让我们选择 1、4、7 和 14.我们将它们混合起来并可视化二进制表示:
Let's take a really simple example. Let's sort four quantities, each of which have 4 binary digits. Let's choose 1, 4, 7, and 14. We'll mix them up and also visualize the binary representation:
Element # 1 2 3 4
Value: 7 14 4 1
Binary: 0111 1110 0100 0001
首先我们将考虑位 0:
First we will consider bit 0:
Element # 1 2 3 4
Value: 7 14 4 1
Binary: 0111 1110 0100 0001
bit 0: 1 0 0 1
现在基数排序算法说我们必须以这样一种方式移动元素(仅考虑位 0)所有零都在左侧，所有零都在右侧.让我们这样做，同时用零位保持元素的顺序，用一位保持元素的顺序.我们可以这样做:
Now the radix sort algorithm says we must move the elements in such a way that (considering only bit 0) all the zeroes are on the left, and all the ones are on the right. Let's do this while preserving the order of the elements with a zero bit and preserving the order of the elements with a one bit. We could do that like this:
Element # 2 3 1 4
Value: 14 4 7 1
Binary: 1110 0100 0111 0001
bit 0: 0 0 1 1
我们的基数排序的第一步已经完成.下一步是考虑下一个(二进制)数字:
The first step of our radix sort is complete. The next step is to consider the next (binary) digit:
Element # 3 2 1 4
Value: 4 14 7 1
Binary: 0100 1110 0111 0001
bit 1: 0 1 1 0
再一次，我们必须移动元素以使相关数字(位 1)按升序排列:
Once again, we must move elements so that the digit in question (bit 1) is arranged in ascending order:
Element # 3 4 2 1
Value: 4 1 14 7
Binary: 0100 0001 1110 0111
bit 1: 0 0 1 1
现在我们必须移动到下一个更高的数字:
Now we must move to the next higher digit:
Element # 3 4 2 1
Value: 4 1 14 7
Binary: 0100 0001 1110 0111
bit 2: 1 0 1 1
然后再次移动它们:
Element # 4 3 2 1
Value: 1 4 14 7
Binary: 0001 0100 1110 0111
bit 2: 0 1 1 1
现在我们移到最后一位(最高位):
Now we move to the last (highest order) digit:
Element # 4 3 2 1
Value: 1 4 14 7
Binary: 0001 0100 1110 0111
bit 3: 0 0 1 0
并做出最后一步:
Element # 4 3 1 2
Value: 1 4 7 14
Binary: 0001 0100 0111 1110
bit 3: 0 0 0 1
现在对值进行了排序.希望这看起来很清楚，但是在到目前为止的描述中，我们已经掩盖了诸如我们如何知道要移动哪些元素?"之类的细节.和我们怎么知道把它们放在哪里?"因此，让我们重复我们的示例，但我们将使用提示中建议的特定方法和顺序来回答这些问题.从第 0 位重新开始:
And the values are now sorted. This hopefully seems clear, but in the description so far we've glossed over the details of things like "how do we know which elements to move?" and "how do we know where to put them?" So let's repeat our example, but we'll use the specific methods and sequence suggested in the prompt, in order to answer these questions. Starting over with bit 0:
Element # 1 2 3 4
Value: 7 14 4 1
Binary: 0111 1110 0100 0001
bit 0: 1 0 0 1
首先让我们构建一个直方图，其中位 0 位置的零位数和位 0 位置的 1 位数:
First let's build a histogram of the number of zero bits in bit 0 position, and the number of 1 bits in bit 0 position:
bit 0: 1 0 0 1
zero bits one bits
--------- --------
1)histogram: 2 2
现在让我们对这些直方图值做一个排他性前缀求和:
Now let's do an exclusive prefix-sum on these histogram values:
zero bits one bits
--------- --------
1)histogram: 2 2
2)prefix sum: 0 2
独占前缀和只是所有前面值的总和.第一个位置没有前面的值，在第二个位置前面的值是 2(第 0 位位置为 0 的元素的数量).现在，作为一个独立的操作，让我们确定每个 0 位在所有 0 位中的相对偏移量，以及每个 1 位在所有 1 位中的相对偏移量:
An exclusive prefix-sum is just the sum of all preceding values. There are no preceding values in the first position, and in the second position the preceding value is 2 (the number of elements with a 0 bit in bit 0 position). Now, as an independent operation, let's determine the relative offset of each 0 bit amongst all the zero bits, and each one bit amongst all the one bits:
bit 0: 1 0 0 1
3)offset: 0 0 1 1
这实际上可以再次使用排他前缀和以编程方式完成，分别考虑 0 组和 1 组，并将每个位置视为其值为 1:
This can actually be done programmatically using exclusive prefix-sums again, considering the 0-group and 1-group separately, and treating each position as if it has a value of 1:
0 bit 0: 1 1
3)ex. psum: 0 1
1 bit 0: 1 1
3)ex. psum: 0 1
现在，给定算法的第 4 步说:
Now, step 4 of the given algorithm says:
4) 结合步骤 2 & 的结果3 确定每个元素的最终输出位置并将其移动到那里
4) Combine the results of steps 2 & 3 to determine the final output location for each element and move it there
这意味着，对于每个元素，我们将选择与其位值(0 或 1)相对应的 histogram-bin 前缀总和值，并添加与其位置相关的偏移量，以确定要移动的位置该元素:
What this means is, for each element, we will select the histogram-bin prefix sum value corresponding to its bit value (0 or 1) and add to that, the offset associated with its position, to determine the location to move that element to:
Element # 1 2 3 4
Value: 7 14 4 1
Binary: 0111 1110 0100 0001
bit 0: 1 0 0 1
hist psum: 2 0 0 2
offset: 0 0 1 1
new index: 2 0 1 3
将每个元素移动到其新索引"位置，我们有:
Moving each element to its "new index" position, we have:
Element # 2 3 1 4
Value: 14 4 7 1
Binary: 0111 1110 0111 0001
根据之前的演练，这正是我们完成第一个数字移动的预期结果.这已经完成了第 1 步，即第一个(最低有效位)数字；我们还有剩余的数字要处理，在每一步创建一个新的直方图和新的前缀和.
Which is exactly the result we expect for the completion of our first digit-move, based on the previous walk-through. This has completed step 1, i.e. the first (least-significant) digit; we still have the remaining digits to process, creating a new histogram and new prefix sums at each step.
注意事项:
- 基数排序，即使在计算机中，也不必严格基于二进制数字.可以用不同大小的数字构建类似的算法，可能由 2、3 或 4 位组成.
- 我们可以对基数排序执行的优化之一是仅根据实际有意义的位数进行排序.例如，如果我们以 32 位值存储数量，但我们知道存在的最大数量是 1023 (2^10-1)，我们不需要对所有 32 位进行排序.在处理前 10 位之后，我们可以停下来，期待正确的排序.
- 这与 GPU 有什么关系?就上述描述而言，不多.实际应用是考虑将并行算法用于直方图、前缀和和数据移动等.基数排序的这种分解允许人们定位和使用已经为这些更基本的操作开发的并行算法，以构建快速并行排序.
- Radix-sort, even in a computer, does not have to be done based strictly on binary digits. It's possible to construct a similar algorithm with digits of different sizes, perhaps consisting of 2,3, or 4 bits.
- One of the optimizations we can perform on a radix sort is to only sort based on the number of digits that are actually meaningful. For example, if we are storing quantities in 32-bit values, but we know that the largest quantity present is 1023 (2^10-1), we need not sort on all 32 bits. We can stop, expecting a proper sort, after proceeding through the first 10 bits.
- What does any of this have to do with GPUs? In so far as the above description goes, not much. The practical application is to consider using parallel algorithms for things like the histogram, the prefix-sums, and the data movement. This decomposition of radix-sort allows one to locate and use parallel algorithms already developed for these more basic operations, in order to construct a fast parallel sort.
以下是一个有效的示例.这可能有助于您理解基数排序.我认为这对您的分配没有帮助，因为此示例在经线级别执行 32 位基数排序，对于单个经线，即.32 个数量.但从理解的角度来看，一个可能的优势是，利用各种 CUDA 内在函数，只需几条指令就可以在扭曲级别完成直方图和前缀和之类的事情.对于您的作业，您将无法使用这些技术，并且您需要提出可以对任意数据集大小进行操作的全功能并行前缀和、直方图等.
What follows is a worked example. This may help with your understanding of radix sort. I don't think it will help with your assignment, because this example performs a 32-bit radix sort at the warp level, for a single warp, ie. for 32 quantities. But a possible advantage from an understanding point of view is that things like histogramming and prefix sums can be done at the warp level in just a few instructions, taking advantage of various CUDA intrinsics. For your assignment, you won't be able to use these techniques, and you will need to come up with full-featured parallel prefix sums, histograms, etc. that can operate on an arbitrary dataset size.
#include <stdio.h>
#include <stdlib.h>
#define WSIZE 32
#define LOOPS 100000
#define UPPER_BIT 31
#define LOWER_BIT 0
__device__ unsigned int ddata[WSIZE];
// naive warp-level bitwise radix sort
__global__ void mykernel(){
__shared__ volatile unsigned int sdata[WSIZE*2];
// load from global into shared variable
sdata[threadIdx.x] = ddata[threadIdx.x];
unsigned int bitmask = 1<<LOWER_BIT;
unsigned int offset = 0;
unsigned int thrmask = 0xFFFFFFFFU << threadIdx.x;
unsigned int mypos;
// for each LSB to MSB
for (int i = LOWER_BIT; i <= UPPER_BIT; i++){
unsigned int mydata = sdata[((WSIZE-1)-threadIdx.x)+offset];
unsigned int mybit = mydata&bitmask;
// get population of ones and zeroes (cc 2.0 ballot)
unsigned int ones = __ballot(mybit); // cc 2.0
unsigned int zeroes = ~ones;
offset ^= WSIZE; // switch ping-pong buffers
// do zeroes, then ones
if (!mybit) // threads with a zero bit
// get my position in ping-pong buffer
mypos = __popc(zeroes&thrmask);
else // threads with a one bit
// get my position in ping-pong buffer
mypos = __popc(zeroes)+__popc(ones&thrmask);
// move to buffer (or use shfl for cc 3.0)
sdata[mypos-1+offset] = mydata;
// repeat for next bit
bitmask <<= 1;
}
// save results to global
ddata[threadIdx.x] = sdata[threadIdx.x+offset];
}
int main(){
unsigned int hdata[WSIZE];
for (int lcount = 0; lcount < LOOPS; lcount++){
unsigned int range = 1U<<UPPER_BIT;
for (int i = 0; i < WSIZE; i++) hdata[i] = rand()%range;
cudaMemcpyToSymbol(ddata, hdata, WSIZE*sizeof(unsigned int));
mykernel<<<1, WSIZE>>>();
cudaMemcpyFromSymbol(hdata, ddata, WSIZE*sizeof(unsigned int));
for (int i = 0; i < WSIZE-1; i++) if (hdata[i] > hdata[i+1]) {printf("sort error at loop %d, hdata[%d] = %d, hdata[%d] = %d
", lcount,i, hdata[i],i+1, hdata[i+1]); return 1;}
// printf("sorted data:
");
//for (int i = 0; i < WSIZE; i++) printf("%u
", hdata[i]);
}
printf("Success!
");
return 0;
}
这篇关于并行基数排序，这个实现实际上是如何工作的?有一些启发式方法吗?的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！
