AN EFFICIENT SORTING ALGORITHM WITH CUDA
이런 방법이 있더군요. 그래서 한번 코드로 구현해보려고 합니다.
이 방법의 개략적인 순서는 다음과 같은 것 같습니다. (영어가 딸려서. --;)
1. 자료값을 크기에 따라서 잘라낸다. (Slice)
2. 잘라낸 것들을 Buckt에 넣는다.
3. Buckt 단위로 정렬한다.
아마도 nVidia GPU의 특성을 활용하기 위해서 이런 방법을 고안한 것 같습니다. global memory에서 그냥 sorting을 하면 편하기는 하겠지만, DDR 램에 썼다 지웠다하면 속도가 느리니, shared memory를 써서 sorting을 해야겠고, shared memory 특성상 다른 블럭과 자료를 공유할 수 없으니 자료 전체를 두고 sorting 할 수가 없어서 일단 전처리로 비슷한 크기끼리 묶는 단계가 필요한 것이겠죠.
그런데 자료 단위가 8백만개 정도는 되야 다른 알고리즘하고 차이가 나고, 그것도 100ms에서 600ms 정도이니. 뭐 그리 큰 필요가 없을 수도 있겠네요. ㅋㅋㅋ
그래도 기왕 배워본 거 연습삼아 조금씩 해볼랍니다.
밑은 코드입니다.
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
|
// AN EFFICIENT SORTING ALGORITHM WITH CUDA
//(http://www.researchgate.net/publication/228848208_An_efficient_sorting_algorithm_with_CUDA/file/9c96052307d2a1add4.pdf) 참조
// Shifu Chen, Jing Qin, Yongming Xie, Junping Zhao, and Pheng-Ann Heng
//Optimizing Parallel Reductiuon 참조
// http://developer.download.nvidia.com/assets/cuda/files/reduction.pdf
#include <iostream>
#include <fstream>
#include <cuda.h>
#include <time.h>
#include <math.h>
using namespace std;
// 테스트 용이므로 일단 자료 크기는 10000으로
// 1D이니까 그냥 블럭사이즈는 512로
// E는 슬라이스의 사이즈
#define DATASIZE 10000
#define BLOCK_SIZE 512
#define EP 256
__global__ void SliceDivision ( float * input, float minValue, unsigned int stepWidth, unsigned int DataSize, unsigned int * SliceOfElements, unsigned int * SliceSizes, unsigned int * OffsetInSlice )
{
//개별 블럭의 좌표
unsigned int tx = threadIdx.x;
//전체 이미지의 좌표
unsigned int x = tx + blockDim.x * blockIdx.x;
// sliceIndex: 어떤 슬라이스에 들어가는 지 필요한 인덱스
// offset: 각 슬라이스 안에서의 offset
unsigned int sliceIndex, offset;
//ceil을 쓰게 되면 sliceIndex[0]에 들어가는 건 가장 작은 숫자밖에 없음. 그래서 floor를 씀
//unsigned int를 쓰기 때문에 음수가 안 나오도록 주의.
if( x < DataSize )
{
sliceIndex = (unsigned int)(floor(( input[x] - minValue ) / stepWidth));
offset = atomicAdd( &SliceSizes[sliceIndex], 1 );
SliceOfElements[x] = sliceIndex;
OffsetInSlice[x] = offset;
}
}
int main()
{
// 테스트에 쓸 숫자 생성
float TestInput[DATASIZE], TestOutput[DATASIZE];
srand(time(NULL));
for( int i = 0; i < DATASIZE; i++ )
{
TestInput[i] = rand() % DATASIZE;
}
//device 설정
float *devInput, *devOutput;
float MaxValue, minValue;
//일단 크기는 아니까
unsigned int MemDataSize = DATASIZE * sizeof(float);
// device 자리 잡아주고
cudaMalloc((void**)&devInput, MemDataSize );
cudaMalloc((void**)&devOutput, MemDataSize );
// 자리 잡았으면 복사
cudaMemcpy( devInput, TestInput, MemDataSize, cudaMemcpyHostToDevice);
// block 크기 설정
// 1D 이니까, 그냥 간단하게...
dim3 dimBlocksize( BLOCK_SIZE );
dim3 dimGridsize( ceil((DATASIZE-1)/(float)BLOCK_SIZE) + 1 );
// 일단 Max값과 min값을 알아내야됨.
// 처음부터 끝까지 휙 둘러보면 되니 이건 CPU에게 맡김.
MaxValue = TestInput[0];
minValue = TestInput[0];
for( int i=1; i < DATASIZE; i++ )
{
if( MaxValue < TestInput[i] )
MaxValue = TestInput[i];
if( minValue > TestInput[i] )
minValue = TestInput[i];
}
// stepWidth 계산
unsigned int stepWidth = (unsigned int)(ceil((MaxValue - minValue) * (float)EP / (float)DATASIZE));
// sliceSizes: 각 슬라이스의 크기를 넣는 array. 전체 크기를 슬라이스의 크기로 나눈 ceil값
unsigned int *SliceSizes, *devSliceSizes;
SliceSizes = new (nothrow) unsigned int[ (unsigned int)( ceil(DATASIZE / (float)stepWidth )) ];
cudaMalloc((void**)&devSliceSizes, ceil(DATASIZE / (float)stepWidth ) * sizeof(unsigned int) );
cudaMemset(devSliceSizes, 0, ceil(DATASIZE / (float)stepWidth ) * sizeof(unsigned int) );
// sliceOfElements: 각 값들의 슬라이스 인덱스 값
unsigned int SliceOfElements[DATASIZE], *devSliceOfElements;
cudaMalloc((void**)&devSliceOfElements, DATASIZE * sizeof(unsigned int) );
cudaMemset(devSliceOfElements, 0, DATASIZE * sizeof(unsigned int) );
// offsetInSlice: 슬라이스 안에서의 offset 값
unsigned int OffsetInSlice[DATASIZE], *devOffsetInSlice;
cudaMalloc((void**)&devOffsetInSlice, DATASIZE * sizeof(unsigned int) );
cudaMemset(devOffsetInSlice, 0, DATASIZE * sizeof(unsigned int) );
// sliceOffsets: 각 값들의 전체 자료에서 offset 값
// unsigned int sliceOffsets[DATASIZE], *devsliceOffsets;
// cudaMalloc((void**)&devsliceOffsets, DATASIZE * sizeof(unsigned int) );
SliceDivision<<< dimGridsize, dimBlocksize >>>( devInput, minValue, stepWidth, DATASIZE, devSliceOfElements, devSliceSizes, devOffsetInSlice );
// 결과물 복사
cudaMemcpy( TestOutput, devOutput, MemDataSize, cudaMemcpyDeviceToHost);
cudaMemcpy( SliceSizes, devSliceSizes, ceil(DATASIZE / (float)stepWidth ) * sizeof(unsigned int), cudaMemcpyDeviceToHost);
cudaMemcpy( OffsetInSlice, devOffsetInSlice, DATASIZE * sizeof(unsigned int), cudaMemcpyDeviceToHost);
cudaMemcpy( SliceOfElements, devSliceOfElements, DATASIZE * sizeof(unsigned int), cudaMemcpyDeviceToHost);
// 결과물 확인
cout << "Max: " << MaxValue << ", min: " << minValue << endl;
for( int i = 0; i < (unsigned int)( ceil(DATASIZE / (float)stepWidth )); i++)
{
cout << i << "th sliceSizes: " << SliceSizes[i] << endl;
}
// 위에 GPU에 마련한 자리 해소. 그때 그때 해놓는 게 편할 듯
cudaFree( devInput );
cudaFree( devOutput );
cudaFree( devSliceSizes );
cudaFree( devSliceOfElements );
cudaFree( devOffsetInSlice );
// cudaFree( devsliceOffsets );
delete SliceSizes;
return 0;
}
|
논문 뒤쪽에도 나와있지만, 이렇게 구현하다보니 각 슬라이스의 크기가 블록크기보다 더 커지면 어떻게 되나 싶더군요. 그래서 다음번에는 슬라이스 중 블록크기보다 더 큰 슬라이스가 있을 때 어떻게 해야할 지 고민 좀 해봐야겠습니다.
댓글 없음:
댓글 쓰기