http://ptjoker95.blogspot.kr/2014/03/cuda-sorting-oddevensorting.html
뭐 알고리즘도 맞고, 결과도 대강맞는 것 같아서 냅두고 다음 정렬로 넘어가려고 잠깐 봤는데, 결과가 이상하게 나오는 겁니다. ㅋㅋㅋ
정렬이 잘 되다가 중간에 잘리고 다른 결과가 나옵니다. 예를 들면,
0 0 0 0 0 1 2 2 2 2 2 ... 495 495 495 496 0 0 0 0 01 1 1 1 1 ...
이런식으로요. 그래서 알고리즘 상에 문제가 있는 줄 알고 뚫어지게 봐도 이상은 없구요. 그래서 혹시나 해서 데이터를 줄여보니, 되더군요. 정확히는 4095 개가 한계이구, 그것보다 커지면 여지없이 잘못된 결과가 나옵니다.
CUDA 강의 중에 작업 순서에 관한 설명이 나오더니 이건가 봅니다.
그러니까 GPU가 돌아갈 때 모든 데이터를 한꺼번에 처리하는 게 아니고, 블럭별로 몇개 블럭씩 (제 경우에는 블럭 사이즈가 512이니까, 4개씩) 처리한다고 하던데, 그래서 문제가 발생한 것 같습니다. ㅋㅋ
물리적 한계라는 게 이런거로군요.
2014/03/27
2014/03/20
CUDA sorting. OddEvenSorting. 병렬 홀짝 정렬
이전에 하던 sorting은 제겐 너무 어렵더군요. --;; 그래서 일단 가장 기초부터 해봤습니다.
이름하야, Odd Even Sorting입니다. 처음에는 홀수(혹은 짝수) 자리에 있는 것과 바로 다음 것을 비교해서 정렬하구요. 그 다음 스텝에는 짝수(혹은 홀수) 자리에 있는 것과 바로 다음 것을 비교 정렬합니다. 그리고 이걸 N(자료크기)번 반복하면 됩니다. 코딩도 매우 간단합니다. 다만, 정렬 시간이 N이 걸리고, GPU에서의 N은 CPU의 N보다 더 오래 걸린다는 점이 단점입니다.
좋은 알고리즘으로하면 log(N)이나 log(N)^2 정도 걸린다는데요. --;; 좋은 알고리즘은 제 머리로는 이해가 어렵다는....
알고리즘은 http://www.dcc.fc.up.pt/~fds/aulas/PPD/1112/sorting.pdf 을 참조했습니다.
코딩이나 개념자체는 쉽습니다.
그리고 참고로 제 PC로 기동했을 때 자료값을 10만개 주면 돌아는 가는데, 이상한 값들이 튀어나오더군요. 문제가 뭔지는 차차 알게되겠죠. 일단은 여기까지.
이름하야, Odd Even Sorting입니다. 처음에는 홀수(혹은 짝수) 자리에 있는 것과 바로 다음 것을 비교해서 정렬하구요. 그 다음 스텝에는 짝수(혹은 홀수) 자리에 있는 것과 바로 다음 것을 비교 정렬합니다. 그리고 이걸 N(자료크기)번 반복하면 됩니다. 코딩도 매우 간단합니다. 다만, 정렬 시간이 N이 걸리고, GPU에서의 N은 CPU의 N보다 더 오래 걸린다는 점이 단점입니다.
좋은 알고리즘으로하면 log(N)이나 log(N)^2 정도 걸린다는데요. --;; 좋은 알고리즘은 제 머리로는 이해가 어렵다는....
알고리즘은 http://www.dcc.fc.up.pt/~fds/aulas/PPD/1112/sorting.pdf 을 참조했습니다.
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
|
#include <iostream>
#include <cuda.h>
#include <time.h>
#include <math.h>
using namespace std;
// 테스트 용이므로 일단 자료 크기는 10000으로
// 1D이니까 그냥 블럭사이즈는 512로
//10만개부터 에러났음. 아마 랜덤 숫자 만들어내는 데, 아니면 GPU메모리 상에서 문제가 발생한 것 같음.
// 만일 화면 데이터를 정렬한다고 하면, 2560x1600 = 4,096,000 픽셀이니까 GPU메모리 상에서의 문제가
// 아니라 랜덤 숫자 만들어내는 곳에서 문제가 발생한 것일 수도...
#define DATASIZE 10000
#define BLOCK_SIZE 512
__global__ void oddevensort ( int * input, int * output, unsigned int len )
{
//개별 블럭의 좌표
unsigned int tx = threadIdx.x;
//전체 이미지의 좌표
unsigned int x = tx + blockDim.x * blockIdx.x;
//이동에 쓸 임시 변수
int temp;
//자료의 길이만큼 돌리는데, 인덱스(i)가 짝수이면 데이터의 짝수자리와 그 다음 숫자를 비교.
//인덱스가 홀수이면 데이터의 홀수자리와 그 다음 숫자를 비교해서 정렬한다.
for( int i=0; i<len; i++)
{
if( i % 2 == 0 )
{
// 길이를 측정안해주면 블럭에 남아있던 자리에 있는 자료가 튀어나올 수 있으니 조심.
if( input[x] > input[x+1] && x < len && x % 2 == 0)
{
temp = input[x+1];
input[x+1] = input[x];
input[x] = temp;
}
}
else
{
if( input[x] > input[x+1] && x < len && x % 2 != 0)
{
temp = input[x+1];
input[x+1] = input[x];
input[x] = temp;
}
}
}
output[x] = input[x];
__syncthreads();
}
int main()
{
// 테스트에 쓸 숫자 생성
int TestInput[DATASIZE], TestOutput[DATASIZE];
srand(time(NULL));
for( int i = 0; i < DATASIZE; i++ )
{
TestInput[i] = rand() % 500;
}
//device 설정
int *devInput, *devOutput;
//일단 크기는 아니까
unsigned int MemDataSize = DATASIZE * sizeof(float);
// device 자리 잡아주고
cudaMalloc((void**)&devInput, MemDataSize );
cudaMalloc((void**)&devOutput, MemDataSize );
cudaMemset( devOutput, 0, MemDataSize );
// 자리 잡았으면 복사
cudaMemcpy( devInput, TestInput, MemDataSize, cudaMemcpyHostToDevice);
// block 크기 설정
// 1D 이니까, 그냥 간단하게...
dim3 dimBlocksize( BLOCK_SIZE );
dim3 dimGridsize( ceil((DATASIZE-1)/(float)BLOCK_SIZE) + 1 );
// 일단 Max값과 min값을 알아내야됨.
// 처음부터 끝까지 휙 둘러보면 되니 이건 CPU에게 맡김.
oddevensort<<< dimGridsize, dimBlocksize >>>( devInput, devOutput, DATASIZE );
// 결과물 복사
cudaMemcpy( TestOutput, devOutput, MemDataSize, cudaMemcpyDeviceToHost);
for( int i=0; i<DATASIZE; i++ )
{
cout << TestOutput[i] << ", ";
if( (i+1)%10 == 0 )
{
cout << endl;
}
}
// 위에 GPU에 마련한 자리 해소. 그때 그때 해놓는 게 편할 듯
cudaFree( devInput );
cudaFree( devOutput );
return 0;
}
|
코딩이나 개념자체는 쉽습니다.
그리고 참고로 제 PC로 기동했을 때 자료값을 10만개 주면 돌아는 가는데, 이상한 값들이 튀어나오더군요. 문제가 뭔지는 차차 알게되겠죠. 일단은 여기까지.
2014/03/10
GPU로 sorting 해보기 step 2
방금 step 2를 올렸는데, 인덱스 상에 안 들어갈 정도로 큰 수를 날려버리는 버그가 있어서 급히 잡았습니다. ㅋㅋ
정렬할 데이터 중에 유난히 큰 수가 있으면 stepWidth가 너무 커지고, 그러다보면 슬라이스 하나의 크기 중에 GPU 블럭보다 큰 게 생길 수 있을 것 같아서, 크기를 조정하는 코드를 새로 넣었습니다.
간단히 설명하면 블럭보다 슬라이스의 크기들이 작아질 때까지 stepWidth를 조정해서 슬라이스의 크기를 조정하게 됩니다. sorting 들어가기 전의 전처리도 꽤나 신경이 많이 쓰이는 군요. ㅋㅋ
기존의 코드에는 없던 SliceLength를 만들어서, 슬라이스들의 총 크기를 알 수 있도록 했습니다. 그래서 GPU 커널에서 슬라이스 인덱스 처리 중에 너무 큰 수가 나올 경우, 가장 마지막 슬라이스에 그 수를 넣도록 했습니다. 어차피 커서 문제일테니 맨 마지막 슬라이스에 넣도록 하는 게 맞겠죠.
for의 조건문을 이용해서, 슬라이스 중 크기가 가장 큰 것이 GPU 블럭보다 작으면 그냥 for문을 패스하도록 했습니다. 처음에는 if를 넣어서 분기했었는데, 굳이 필요없겠더군요.
안전을 위해서 stepWidth가 50보다 작아지면 프로그램을 종료하도록 했습니다. 최악의 경우에는 슬라이스 크기가 1나 2개 정도로 될 수도 있겠더군요.
다음에는 나뉘어진 Slice를 Bucket에 넣는 작업입니다. Bucket의 크기는 GPU의 블럭 크기와 동일하며, 하나의 Bucket에 크기를 넘지 않는 범위내에서 가장 많은 Slice를 넣어야 좀 더 효율적일 것 같습니다.
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
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
|
// 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;
// 테스트 용이므로 일단 자료 크기는 1000으로
// 1D이니까 그냥 블럭사이즈는 512로
// EP는 슬라이스의 사이즈
#define DATASIZE 1000
#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 SliceLength )
{
//개별 블럭의 좌표
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));
//sliceIndex가 슬라이스들의 크기보다 크면 가장 큰 곳에 넣음
if( sliceIndex > SliceLength-1 )
sliceIndex = SliceLength-1;
offset = atomicAdd( &SliceSizes[sliceIndex], 1 );
SliceOfElements[x] = sliceIndex;
OffsetInSlice[x] = offset;
}
}
int main()
{
// 테스트에 쓸 숫자 생성
float TestInput[DATASIZE], TestOutput[DATASIZE];
unsigned int MaxSliceSize, SliceLength;
srand(time(NULL));
// slice 크기 조정하는 코드 확인하기 위해서 맨 마지막 자리에 꽤 큰 숫자를 넣어보겠습니다.
for( int i = 0; i < DATASIZE-2; i++ )
{
TestInput[i] = rand() % DATASIZE;
}
TestInput[DATASIZE-1] = 10000;
//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;
SliceLength = (unsigned int)( ceil(DATASIZE / (float)stepWidth ));
cout << "SliceLength: " << SliceLength << endl;
SliceSizes = new (nothrow) unsigned int[ SliceLength ];
cudaMalloc((void**)&devSliceSizes, (unsigned int)(SliceLength * sizeof(unsigned int)) );
cudaMemset(devSliceSizes, 0, SliceLength * 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, SliceLength );
// 결과물 복사
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);
MaxSliceSize = SliceSizes[0];
for( int i = 1; i < (unsigned int)( ceil(DATASIZE / (float)stepWidth )); i++)
{
//슬라이스 중에 가장 큰 크기를 MaxSliceSize에 넣음
if( MaxSliceSize < SliceSizes[i] )
{
MaxSliceSize = SliceSizes[i];
}
}
//슬라이스 중에 가장 큰 크기가 BLOCK_SIZE보다 클 경우에는 stepWidth를 조정해서 다시 계산.
// 이 때, sliceSize를 가지는 array의 크기가 달라질 테니 그것도 같이 조정
// 일단 한번 크면 또 클 수도 있으니 루프를 써서 일정 크기로 줄어들게 함. 일단은 64개 씩
//크기가 1보다 같거나 작을 때까지 돌림. 너무 작아지면 일단은 프로그램 정지
for( unsigned int i = 1; MaxSliceSize > 512; i++ )
{
cout << "Max: " << MaxValue << ", min: " << minValue << ", stepWidth: " << stepWidth << ", MaxSliceSize: " << MaxSliceSize << ", i:" << i << endl;
//stepWidth 조정
stepWidth = (unsigned int)(ceil((MaxValue - minValue) * (float)EP / ((float)(DATASIZE * 2 * i ))));
SliceLength = (unsigned int)( ceil(DATASIZE / (float)stepWidth ));
cout << "SliceLength: " << SliceLength << endl;
// 기존의 슬라이스 크기를 넣은 array 삭제
delete[] SliceSizes;
// 슬라이스 크기를 넣는 array 다시 만들기
SliceSizes = new (nothrow) unsigned int[ SliceLength ];
// GPU 메모리도 삭제
cudaFree( devSliceSizes );
// GPU 메모리도 다시 만들고
cudaMalloc((void**)&devSliceSizes, (unsigned int)(SliceLength * sizeof(unsigned int)) );
// GPU 메모리 청소
cudaMemset(devSliceSizes, 0, (unsigned int)(SliceLength * sizeof(unsigned int)) );
//커널 다시 실행
SliceDivision<<< dimGridsize, dimBlocksize >>>( devInput, minValue, stepWidth, DATASIZE, devSliceOfElements, devSliceSizes, devOffsetInSlice, SliceLength );
// 결과물 복사
cudaMemcpy( TestOutput, devOutput, MemDataSize, cudaMemcpyDeviceToHost);
cudaMemcpy( SliceSizes, devSliceSizes, (unsigned int)(ceil(DATASIZE / (float)stepWidth ) * sizeof(unsigned int)), cudaMemcpyDeviceToHost);
cudaMemcpy( OffsetInSlice, devOffsetInSlice, DATASIZE * sizeof(unsigned int), cudaMemcpyDeviceToHost);
cudaMemcpy( SliceOfElements, devSliceOfElements, DATASIZE * sizeof(unsigned int), cudaMemcpyDeviceToHost);
// 슬라이스 크기가 BLOCK_SIZE보다 작은지 다시 확인
MaxSliceSize = SliceSizes[0];
for( int i = 1; i < (unsigned int)( ceil(DATASIZE / (float)stepWidth )); i++)
{
//슬라이스 중에 가장 큰 크기를 MaxSliceSize에 넣음
if( MaxSliceSize < SliceSizes[i] )
{
MaxSliceSize = SliceSizes[i];
}
}
//stepWidth 가 50보다 작으면 뭔가가 문제있는 걸로 판단하고 일단 계산 스톱. 잘못하면 무한 루프 걸림
if( stepWidth < 50 )
return 1;
}
cout << "Max: " << MaxValue << ", min: " << minValue << ", stepWidth: " << stepWidth << ", MaxSliceSize: " << MaxSliceSize << endl;
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;
}
|
정렬할 데이터 중에 유난히 큰 수가 있으면 stepWidth가 너무 커지고, 그러다보면 슬라이스 하나의 크기 중에 GPU 블럭보다 큰 게 생길 수 있을 것 같아서, 크기를 조정하는 코드를 새로 넣었습니다.
간단히 설명하면 블럭보다 슬라이스의 크기들이 작아질 때까지 stepWidth를 조정해서 슬라이스의 크기를 조정하게 됩니다. sorting 들어가기 전의 전처리도 꽤나 신경이 많이 쓰이는 군요. ㅋㅋ
기존의 코드에는 없던 SliceLength를 만들어서, 슬라이스들의 총 크기를 알 수 있도록 했습니다. 그래서 GPU 커널에서 슬라이스 인덱스 처리 중에 너무 큰 수가 나올 경우, 가장 마지막 슬라이스에 그 수를 넣도록 했습니다. 어차피 커서 문제일테니 맨 마지막 슬라이스에 넣도록 하는 게 맞겠죠.
for의 조건문을 이용해서, 슬라이스 중 크기가 가장 큰 것이 GPU 블럭보다 작으면 그냥 for문을 패스하도록 했습니다. 처음에는 if를 넣어서 분기했었는데, 굳이 필요없겠더군요.
안전을 위해서 stepWidth가 50보다 작아지면 프로그램을 종료하도록 했습니다. 최악의 경우에는 슬라이스 크기가 1나 2개 정도로 될 수도 있겠더군요.
다음에는 나뉘어진 Slice를 Bucket에 넣는 작업입니다. Bucket의 크기는 GPU의 블럭 크기와 동일하며, 하나의 Bucket에 크기를 넘지 않는 범위내에서 가장 많은 Slice를 넣어야 좀 더 효율적일 것 같습니다.
2014/03/04
GPU로 sorting 해보기 step 1
뭐니뭐니해도 GPU라면 매트릭스 계산과 sorting이 되야겠죠. 그래서 sorting 알고리즘을 알아보니,
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 정도이니. 뭐 그리 큰 필요가 없을 수도 있겠네요. ㅋㅋㅋ
그래도 기왕 배워본 거 연습삼아 조금씩 해볼랍니다.
밑은 코드입니다.
논문 뒤쪽에도 나와있지만, 이렇게 구현하다보니 각 슬라이스의 크기가 블록크기보다 더 커지면 어떻게 되나 싶더군요. 그래서 다음번에는 슬라이스 중 블록크기보다 더 큰 슬라이스가 있을 때 어떻게 해야할 지 고민 좀 해봐야겠습니다.
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;
}
|
논문 뒤쪽에도 나와있지만, 이렇게 구현하다보니 각 슬라이스의 크기가 블록크기보다 더 커지면 어떻게 되나 싶더군요. 그래서 다음번에는 슬라이스 중 블록크기보다 더 큰 슬라이스가 있을 때 어떻게 해야할 지 고민 좀 해봐야겠습니다.
피드 구독하기:
글 (Atom)