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
|
#include <iostream>
#include <CImg.h>
#include <cuda.h>
using namespace cimg_library;
using namespace std;
//블럭 사이즈는 16x16 = 256. 블럭 하나의 한계치가 얼마인지 모르지만, 16x16이 안전할 듯
//한계치는 천천히 배워가면 됨.
//depth와 spectrum이 추가되어서 혹시나해서 12로 바꿈
#define BLOCK_SIZE 16
__global__ void add( unsigned char * input, unsigned char * output, unsigned int W, unsigned int H, unsigned int D, unsigned int S )
{
//개별 블럭의 좌표
unsigned int tx = threadIdx.x;
unsigned int ty = threadIdx.y;
//전체 이미지의 좌표
unsigned int x = tx + blockDim.x * blockIdx.x;
unsigned int y = ty + blockDim.y * blockIdx.y;
// 흑백으로 할 것이기 때문에, 높이는 그냥 H로 지정
if( x < W && y < H )
{
//원본이미지 RGB값을 조정해서 복사할 곳에 넣어야 함
//R부터 시작하고, 각 픽셀에 해당하는 G와 B는 높이의 2배, 3배 위치에 있으므로 y에 H를 더해서 조정함
output[ x + W * y ] = (unsigned char)(input[ x + W * y ] * 0.2126 + input[ x + W * (y + H) ] * 0.7152 + input[ x + W * ( y + 2 * H) ] * 0.0722);
__syncthreads();
}
}
int main()
{
// GPU 메모리
// CImg가 데이터를 unsigned char로 하기 때문에 GPU에 쓸 메모리도 unsigned char로 해야됨.
// 그래야 픽셀 단위로 처리가 가능해짐
unsigned char *devInputImage;
unsigned char *devOutputImage;
// 읽어온 이미지의 넓이과 높이, 색깔 채널을 넣을 장소
unsigned int ImageWidth, ImageHeight, ImageDepth, ImageSpectrum;
unsigned int ImageSize;
bool ImageShared;
// 이미지 읽어오기
CImg<unsigned char> image("images.jpg");
//크기 및 기타 등등
ImageWidth = image._width;
ImageHeight = image._height;
ImageDepth = image._depth;
ImageSpectrum = image._spectrum;
ImageShared = image._is_shared;
//복사할 자리 만들기
//흑백으로 할 것이기 때문에 spectrum을 1로 지정
CImg<unsigned char> visu( ImageWidth, ImageHeight, ImageDepth, 1, ImageShared);
// CImg에서는 이미지의 크기를 .size()로 간단하게 알 수 있음
ImageSize = image.size();
// GPU에 자리 마련하기, 크기는 이미지 크기를 그대로 적으면 됨.
// unsigned char의 크기가 1이기 때문에 생략. 넣어도 상관없을 듯.
cudaMalloc((void**)&devInputImage, ImageSize );
// 복사할 자리는 원래 크기의 1/3이기 때문에, 이미지의 넓이과 높이만으로 지정
cudaMalloc((void**)&devOutputImage, ImageWidth * ImageHeight );
// 노파심에 GPU 아웃풋 자리를 일단 청소.
cudaMemset(devOutputImage, 0, ImageSize );
// 자리도 마련했으니 데이터를 복사해줌.
cudaMemcpy( devInputImage, image, ImageSize, cudaMemcpyHostToDevice);
// 블럭의 크기 지정
// 이때 Grid의 크기에 주의해야됨.
// CImg의 데이터는 unsigned char의 크기로, RRRRRR...GGGGGG...BBBBB..., 이런식으로 저장됨.
// 즉, 이미지 크기가 225라면, 넓이는 225 그대로 해줘도 되지만, 높이는 225 x 3으로 해줘야 데이터가 잘리지 않음
dim3 dimBlocksize( BLOCK_SIZE, BLOCK_SIZE );
dim3 dimGridsize( ceil((ImageWidth-1)/BLOCK_SIZE) + 1, (ceil((ImageHeight-1)/BLOCK_SIZE) + 1) * 3 );
// GPU kernel 실행
add<<< dimGridsize, dimBlocksize >>>( devInputImage, devOutputImage, ImageWidth, ImageHeight, ImageDepth, ImageSpectrum );
//일이 끝났으니 결과물을 CPU로 복사함.
//결과물을 GPU에서 CPU로 복사. 복사할 때에도 크기지정은 잘 해줘야 됨.
cudaMemcpy(visu, devOutputImage, ImageWidth * ImageHeight, cudaMemcpyDeviceToHost);
//복사했으니 제대로 되었는지 확인
CImgDisplay main_disp(image,"Original"), copy_disp(visu,"copy");
while( !main_disp.is_closed() && !copy_disp.is_closed() )
main_disp.wait();
// 위에 GPU에 마련한 자리 해소. 그때 그때 해놓는 게 편할 듯
cudaFree( devInputImage );
cudaFree( devOutputImage );
return 0;
}
|
중요한 건,
- 흑백이미지는 RGB보다 크기가 1/3
이라는 점입니다.
각각의 RGB값을 계산한 것을 하나의 흑백이미지 저장소에 잘 집어넣어야 문제가 안 생길 것 같습니다. 그럼 이만.
댓글 없음:
댓글 쓰기