프로그래밍/CUDA에 해당되는 글 3건

Image Processing with CUDA

프로그래밍/CUDA 2015. 7. 28. 02:39

  

    Cuda_ex1.zip



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
// main.cpp
 
#include <cv.hpp>
#include "kernel.h"
using namespace cv;
 
int main()
{
    // Open a webcamera
    VideoCapture camera(0);
    Mat frame;
    if (!camera.isOpened()) 
        return -1;
    
    camera >> frame;
 
    // create CPU/GPU shared images - one for the initial and one for the result
    Mat sGray(frame.size(),CV_8U,createImageBuffer(frame.size().width * frame.size().height));
    Mat dGray(frame.size(),CV_8U,createImageBuffer(frame.size().width * frame.size().height));
    Mat eGray(frame.size(),CV_8U,createImageBuffer(frame.size().width * frame.size().height));
    cvtColor(frame, dGray, CV_BGR2GRAY);
    cvtColor(frame, eGray, CV_BGR2GRAY);
 
    // Create the capture windows
    namedWindow("Source");
    namedWindow("Greyscale");
    namedWindow("Blurred");
    namedWindow("Sobel");
    
    char c;
 
    // Loop while capturing images
    while (1)
    {
        // Capture the image and store a gray conversion for the gpu
        camera >> frame;
        cv::cvtColor(frame, sGray, CV_BGR2GRAY);
        
        boxfilter(frame.size().width, frame.size().height, sGray.data, dGray.data, 3, 3);
        //boxfilterCPU(frame.size().width, frame.size().height, sGray.data, dGray.data, 3, 3);
        sobelfilter(frame.size().width, frame.size().height, dGray.data, eGray.data);
 
        // Show the results
        cv::imshow("Source", frame);
        cv::imshow("Greyscale", sGray);
        cv::imshow("Blurred", dGray);
        cv::imshow("Sobel", eGray);
 
        c = cv::waitKey(10);
        
        if (c == 27)
            break;
    }
 
    // Exit
    destroyImageBuffer(sGray.data);
    destroyImageBuffer(dGray.data);
    destroyImageBuffer(eGray.data);
 
    return 0;
}
 
cs


1
2
3
4
5
6
7
8
9
10
11
12
13
14
// kernel.h
 
#ifndef _KERNEL_H_
#define _KERNEL_H_
 
#include <iostream>
 
void boxfilter(int iw, int ih, unsigned char *source, unsigned char *dest, int bw, int bh);
void boxfilterCPU(int iw, int ih, unsigned char *src, unsigned char *dst, int bw, int bh);
void sobelfilter(int iw, int ih, unsigned char *source, unsigned char *dest);
 
unsigned char* createImageBuffer(unsigned int bytes);
void           destroyImageBuffer(unsigned char* bytes);
#endif
cs

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
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include "kernel.h"
 
void boxfilterCPU(int iw, int ih, unsigned char *src, unsigned char *dst, int bw, int bh)
{
    for (int m = 0; m < ih; m++)
    {
        for (int n = 0; n < iw; n++)
        {
            int count = 0;
            float sum = 0.0;
 
            for (int j = -(bh / 2); j <= (bh / 2); j++)
            {
                for (int i = -(bw / 2); i <= (bw / 2); i++)
                {
                    // Verify that this offset is within the image boundaries
                    if ((n + i) < iw && (n + i) >= 0 && (m + j) < ih && (m + j) >= 0)
                    {
                        sum += (float)src[((m + j) * iw) + (n + i)];
                        count++;
                    }
                }
            }
 
            // Average the sum
            sum /= (float)count;
            dst[(m * iw) + n] = (unsigned char)sum;
        }
    }
}
 
__global__ void boxfilter_kernel(int iw, int ih, unsigned char *source, unsigned char *dest, int bw, int bh)
{
    // Calculate our pixel's location
    int x = (blockIdx.x * blockDim.x) + threadIdx.x;
    int y = (blockIdx.y * blockDim.y) + threadIdx.y;
 
    // Variables to store the sum
    int count = 0;
    float sum = 0.0;
 
    // Do the blur operation by summing the surround pixels
    for (int j = -(bh / 2); j <= (bh / 2); j++)
    {
        for (int i = -(bw / 2); i <= (bw / 2); i++)
        {
            // Verify that this offset is within the image boundaries
            if ((x + i) < iw && (x + i) >= 0 && (y + j) < ih && (y + j) >= 0)
            {
                sum += (float)source[((y + j) * iw) + (x + i)];
                count++;
            }
        }
    }
 
    // Average the sum
    sum /= (float)count;
    dest[(y * iw) + x] = (unsigned char)sum;
}
 
__global__ void sobelfilter_kernel(int iw, int ih, unsigned char *source, unsigned char *dest)
{
    // Calculate our pixel's location
    int x = (blockIdx.x * blockDim.x) + threadIdx.x;
    int y = (blockIdx.y * blockDim.y) + threadIdx.y;
 
    // Operate only if we are in the correct boundaries
    if (x > 0 && x < iw - 1 && y > 0 && y < ih - 1)
    {
        int gx = -source[iw*(y - 1) + (x - 1)] + source[iw*(y - 1) + (x + 1)] +
            -2 * source[iw*(y)+(x - 1)] + 2 * source[iw*(y)+(x + 1)] +
            -source[iw*(y + 1) + (x - 1)] + source[iw*(y + 1) + (x + 1)];
        int gy = -source[iw*(y - 1) + (x - 1)] - 2 * source[iw*(y - 1) + (x)]
            - source[iw*(y - 1) + (x + 1)] +
            source[iw*(y + 1) + (x - 1)] + 2 * source[iw*(y + 1) + (x)] +
            source[iw*(y + 1) + (x + 1)];
        dest[iw*y + x] = (int)sqrt((float)(gx)*(float)(gx)+(float)(gy)*(float)(gy));
    }
}
 
 
 
 
void boxfilter(int iw, int ih, unsigned char *source, unsigned char *dest, int bw, int bh)
{
    // allocate memory for the bitmap in GPU memory
    unsigned char *dev_source, *dev_dest;
    cudaHostGetDevicePointer(&dev_source, source, 0);
    cudaHostGetDevicePointer(&dev_dest, dest, 0);
 
    //cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags()
 
    // Run the boxfilter kernel
    dim3 blocks(iw / 16, ih / 16);
    dim3 threads(16, 16);
    
    // Execute the kernel
    boxfilter_kernel << <blocks, threads >> >(iw, ih, dev_source, dev_dest, bw, bh);
    cudaThreadSynchronize();
}
 
void sobelfilter(int iw, int ih, unsigned char *source, unsigned char *dest)
{
    // allocate memory for the bitmap in GPU memory
    unsigned char *dev_source, *dev_dest;
    cudaHostGetDevicePointer(&dev_source, source, 0);
    cudaHostGetDevicePointer(&dev_dest, dest, 0);
 
    // Run the boxfilter kernel
    dim3 blocks(iw / 16, ih / 16);
    dim3 threads(16, 16);
 
    // Execute the kernel
    sobelfilter_kernel << <blocks, threads >> >(iw, ih, dev_source, dev_dest);
    cudaThreadSynchronize();
}
 
unsigned char* createImageBuffer(unsigned int bytes)
{
    unsigned char *ptr = NULL;
    cudaSetDeviceFlags(cudaDeviceMapHost);
    cudaHostAlloc(&ptr, bytes, cudaHostAllocMapped);
    return ptr;
}
 
void destroyImageBuffer(unsigned char* bytes)
{
    cudaFreeHost(bytes);
}
cs



저작자표시

'프로그래밍 > CUDA' 카테고리의 다른 글

CUDA 스레드 구조  (0) 2014.12.06
CUDA 기본 문법 정리  (0) 2014.12.06
,

CUDA 스레드 구조

프로그래밍/CUDA 2014. 12. 6. 14:13

CUDA의 스레드는 계층 구조로 이루어져있다. 스레드-블록-그리드로 이루어져 있다. 스레드가 모여 블록을

이루고 블록이 모여 그리드를 이루게 된다.



위 사진과 같이 스레드가 모여 블록을 이루고 블록이 모여 그리드를 이룬다.


CUDA의 블록은 스레드가 모인 집합이다. 하나의 블록은 1~512개의 스레드를 가질 수 있다.


'프로그래밍 > CUDA' 카테고리의 다른 글

Image Processing with CUDA  (0) 2015.07.28
CUDA 기본 문법 정리  (0) 2014.12.06
,

CUDA 기본 문법 정리

프로그래밍/CUDA 2014. 12. 6. 13:39

<CUDA C언어 기본 문법 정리>


__global__

· 디바이스에서 실행된다. 호스트에서 호출할 수 있지만, 디바이스에서 호출할 수 없다.

· 디바이스로 실행하는 커널 함수 지정에 사용할 수 있다.


1
2
3
__global__ function<<< >>> () 
 
{ ...}


* 주의사항

1. 리턴값은 항상 void이다.(리턴값을 지정할 수 없다. (x))

2. <<<, >>>를 이용하여 실행 시 블록과 스레드 지정이 가능하다.

3. 재귀 호출은 불가능

4. 함수내 static변수를 가질 수 없다.

5. 가변형 인수를 가질 수 없음( __global__ function<<< ... >>> (int a, ...) (x)

6. __global__로 지정한 함수의 포인터를 이용할 수 있다.

7. __host__와 동시에 이용할 수 없다.

8. 디바이스에서 처리가 완료되기 전에 호출한 즉시 반환하여 비동기 동작

9. 공유 메모리를 이용하요 256바이트까지의 인수 사용이 가능하다.



__device__

· 디바이스에서 실행된다. 디바이스에서 호출할 수 있고, 호스트에서 호출이 불가능하다. 디바이스

코드 중에 작성하여 디바이스 내에서의 실행되는 서브함수로 사용한다.


1
2
__device__ int function(int a, int b)
{...}


* 주의사항

1. 재귀호출 사용 불가

2. 함수내 static변수를 가질 수 없다.

3. 가변형 인수를 가질 수 없다.

4. __device__로 지정한 함수의 포인터는 사용할 수 없다.



__host__ 

· 호스트에서 실행된다. 호스트에서 호출할 수 있고, 디바이스에서 호출할 수 없다. 호스트에서 보통 사용하는 함수가 된다.


1
2
__host__ int function(int a, int b)
{ ... }


* 주의사항

1. __host__, __global__, __device__가 지정되지 않은 경우, __host__를 지정한 것과 동일함

2. __global__과 동시에 사용할 수 없다.

3. __device__와 동시에 사용하여, 호스트와 디바이스 양쪽에서 사용할 수 있는 함수로 작성할 수 있다.




<변수 수식어>
1. __device__
· 글로벌 메모리 영역에 할당되어 프로그램이 종료될 때까지 유효하다. 모든 스레드가 액세스 할 수 있고, 호스트 측에서는 API함수를 통해서 읽기와 쓰기가 가능하다.

2. __constant__

· 상수 메모리(constant memory) 영역에 할당되어 프로그램이 종료될 때까지 유효하다. 모든 스레드가 액세스하며 읽기만 가능하다. 호스트에서 cudaMemoryToSymbol()의 API를 통해서 값을 쓸 수 있다. 상수 캐시(constant cache)가 함께 사용된다.

3. __shared__

· 공유 메모리 영역에 할당되어 실행 중인 스레드 블록 상에서 유효하다. 블록 내의 스레드는 액세스하여 읽고 쓰기가 가능하다.


'프로그래밍 > CUDA' 카테고리의 다른 글

Image Processing with CUDA  (0) 2015.07.28
CUDA 스레드 구조  (0) 2014.12.06
,
1

티스토리툴바