기본적인 설명은, CImg는 자료를 unsigned char로 만들어 놓는데, 그 순서는 RRRRRR....GGGGG.....BBBBB...., 이런 식으로 됩니다. (참조: http://cimg.sourceforge.net/reference/group__cimg__storage.html)
그런데 이미지 조작이라도 하려면 각각의 값을 따로 조작할 수 있어야 하는데, 그게 안되더라는 거죠. 한참 삽질을 한 다음에 이유를 알게되었습니다.
우선 코드...
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
|
#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
// 20140220 복사는 되는데, 자료값이 엉망이 됨. ㅋㅋㅋ
__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;
// 일단 모든 RGB를 모두 포함하기 위해서, 기본 이미지의 높이(H)보다 3배 더 크게 잡아줌
if( x < W && y < 3 * H )
{
//원본이미지를 아웃풋에 복사함
//RGB에서 G를 없애는 작업
if( H <= y && y < 2 * H )
output[ x + W * y ] = 0;
else
output[ x + W * y ] = input[ x + W * y ];
//나머지는 그대로 놔둠.
__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;
//복사할 자리 만들기
CImg<unsigned char> visu( ImageWidth, ImageHeight, ImageDepth, ImageSpectrum, ImageShared);
// CImg에서는 이미지의 크기를 .size()로 간단하게 알 수 있음
ImageSize = image.size();
// GPU에 자리 마련하기, 크기는 이미지 크기를 그대로 적으면 됨.
// unsigned char의 크기가 1이기 때문에 생략. 넣어도 상관없을 듯.
cudaMalloc((void**)&devInputImage, ImageSize );
cudaMalloc((void**)&devOutputImage, ImageSize );
// 노파심에 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, ImageSize, 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;
}
|
처음에는 GPU에서 쓸 메모리 값을 float으로 했더랬습니다. 그렇게 하니, 픽셀의 (x,y)위치가 마음먹은 대로 컨트롤이 안되는 겁니다. 한참을 이래저래 해보니 이유를 알겠더군요.
CImg는 자료를 unsigned char로 저장합니다. 그리고 그 크기는 float 크기의 1/4입니다. 그러니까, GPU로 가져간 자료는 4개의 픽셀을 한꺼번에 들고 갔던거죠. 그러니 GPU에서 thread의 x,y자표가 엉망이 되는거죠. ㅋ
그래서 GPU에서 unsigned char로 바꾸니, 이젠 일부 자료가 날아가는 겁니다. 즉, Grid값이 잘못 된거죠. 그래서 Grid값을 3배 해주니 되더군요.
간단하게 설명하면, 이미지가 3개 위아래로 나란히 배열되어 있다고 생각하면 될 것 같습니다. 맨 위에는 R이미지, 중간에는 G이미지, 맨밑이 B이미지인거죠.
그렇게 해서 GPU의 y값을 이미지 높이로 조정할 수 있게 되었습니다.
다음 번에는 gray 로 변환하는 걸 해볼랍니다.
댓글 없음:
댓글 쓰기