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를 넣어야 좀 더 효율적일 것 같습니다.
댓글 없음:
댓글 쓰기