이름하야, 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만개 주면 돌아는 가는데, 이상한 값들이 튀어나오더군요. 문제가 뭔지는 차차 알게되겠죠. 일단은 여기까지.
댓글 없음:
댓글 쓰기