프로그래밍에 해당되는 글 41건

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
,

ubuntu 14.04 LTS iTorch 설치

프로그래밍 2015. 4. 18. 03:56

iTorch 설치 가이드


다양한 포멧 출력 및 시각 효과(그래프) 등을 지원하는 iTorch 설치 방법을 알아보겠습니다.

iTorch 는 사진, 오디오, 비디오 등 다양한 형식 출력을 지원합니다.


사진


동영상


그래프


itorch를 설치하긴 위해선 연관 라이브러리를 설치해야 합니다.



  https://github.com/facebook/iTorch


에서 설치 방법을 보면 Torch-7를 설치해야 합니다.


Torch설치는 지난 포스팅을 참조하면 됩니다.


그리고 요구조건인 IPython을 설치해야 하는데 설치 방법은 다음과 같습니다.


1. IPython 설치


1.1 pip 명령어로 ipython을 설치

 $ pip install ipython


1.2 ipython-notebook을 설치

 $ pip install ipython-notebook


를 설치하면 설치가 완료됩니다.

 $ ipython --version


를 입력해서 만약 버전이 2.2 보다 낮으면 업데이트를 해야합니다.

업데이트를 하려면 다음과 같이 입력하면 됩니다.


 $ pip install --upgrade ipython[all]


업데이를 완료한 뒤 ipython --version를 입력하면 아래 사진과 같이 버전이 업데이트 된것을 볼 수 있습니다.



2. iTorch 설치


2.1 Dependencies 라이브러리를 다운받습니다.

 $ sudo apt-get install libzmq3-dev


2.2 git허브에서 소스를 다운받습니다.

 $ git clone https://github.com/facebook/iTorch.git


2.3 다운받은 폴더로 들어가 make합니다.

 $ cd iTorch


 $ luarocks make


만약 sudo 관리자로 설치했을 경우(혹은 global로 설치했을 경우)

다음과 같이 입력합니다.


 $ sudo env "PATH=$PATH" luarocks make

 $ sudo chown -R $USER $(dirname $(ipython locate profile))


이제 설치가 완료되었습니다.


2.4 실행

 $ itorch notebook # notebook mode

   or

 $ itorch             # consol mode
   or
 $ itorch qtconsole # Qt mode


2.5 예제

iTorch notebook를 이용해서 Demo를 테스트 해볼수 있습니다.(아래 주소 참조)

http://nbviewer.ipython.org/github/facebook/iTorch/blob/master/iTorch_Demo.ipynb


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

ubuntu 14.04 LTS Torch 설치  (0) 2015.04.17
Lua 프로그래밍  (0) 2015.04.16
안드로이드 가상 머신 지니모션(Genymotion) 설치  (0) 2015.04.09
멀티 스레드  (0) 2014.12.13
파이썬 argparse  (0) 2014.11.08
,

ubuntu 14.04 LTS Torch 설치

프로그래밍 2015. 4. 17. 15:45

1. 터미널 실행


2. 우분투 root 비밀번호 설정

 $ sudo passwd root


3. 관리자 로그인

 $ su

   암호 : ___


4. torch 및 관련 패키지 다운로드

 $ curl -s https://raw.githubusercontent.com/torch/ezinstall/master/install-all | bash


5. 실행

단순히 실행하려면

 $ luajit -ltorch

을 입력하면 아래와 같이 shell이 실행된다.



5.1 th interpreter 설정

 $ th -lparallel -loptim -lpl -limage

 $ th


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

ubuntu 14.04 LTS iTorch 설치  (2) 2015.04.18
Lua 프로그래밍  (0) 2015.04.16
안드로이드 가상 머신 지니모션(Genymotion) 설치  (0) 2015.04.09
멀티 스레드  (0) 2014.12.13
파이썬 argparse  (0) 2014.11.08
,

Lua 프로그래밍

프로그래밍 2015. 4. 16. 02:12


1. 변수 선언

1
2
3
4
5
6
7
num = 42 (숫자, 모든 숫자는 더블형으로 저장된다)
 
s = 'hello world' (문자열)
 
u = "double-quotes are also fine" (문자열)
 
t  =  nil (정의되지 않은 t 선언) 
cs


2. 반복문

1
2
3
4
5
while num < 50 do
 
    num = num + 1 (++, +=와 같음)
 
end
cs

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
//---------------------------------------------------------------
// 상승 반복문(1씩 증가)
 
upSum = 0
 
for i = 1, 100 do
   upSum = upSum + i
end
 
//---------------------------------------------------------------
// 하강 반복문(1씩 감소)
 
downSum = 0
 
for j = 100, 1, -1 do
   downSum = downSum + j
end
cs

1
2
3
4
repeat
   print('down')
   num = num - 1
until num == 0
cs


3. 조건문

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
if num > 40 then
 
    print('over 40')
 
elseif s ~= 'walternate' then (~= 는 같지 않으면을 의미)
 
    io.write('not over 40\n') (stdout)
 
else
 
    thisIsGlobal = 5
 
    local line = io.read (stdin, line 읽기)
 
    print('Hello, ' .. line)
 
end
 
//--------------------------------------------------------------------------------------//
 
aBoolValue = false
 
if not aBoolValue then 
   print('twas false') 
end
 
//--------------------------------------------------------------------------------------//
 
ans = aBoolValue and 'yes' or 'no'  -->  no
 
cs


4. 함수

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
//--------------------------------------------------------------------------------------//
// fib함수는 변수로 숫자를 받아 피보나치 수열 출력
 
function fib(n)
   if n < 2 then return 1 end
   return fib(n - 2) + fib(n - 1)
end
 
//--------------------------------------------------------------------------------------//
// 함수 안에 함수를 생성 가능
 
function adder(x)
   return function(y) return x + y end
end
 
a1 = adder(9)
a2 = adder(36)
print(a1(16)) -> 결과 = 25출력
print(a2(64)) -> 결과 = 100출력
cs

1
2
3
4
5
6
7
8
9
10
x, y, z = 1, 2, 3, 4
 
// a, b, c를 입력받아 6개 숫자를 리턴하는 함수 선언
function bar(a, b, c)
     print(a, b, c)
     return 4, 8, 15, 16, 23, 42
end
 
x, y = bar('zaphod')  -> prints "zaphod nil nil"
 
cs


5. 테이블

1
2
3
4
// 사전(dictionary) 생성
t = {key1 = "value1", key2 = false}
 
print(t.key1) -> 'value1'출력
cs

1
2
3
4
5
v = {"value1", "value2", 1.21, "endList"}
 
for i = 1, #v do (#v는 리스트 사이즈)
   print(v[i])
end
cs


6. 기타


 PDF 파일 다운로드

 luarefv51.pdf

   


좀더 자세한 내용은

string library
table library
math library
io library
os library

참조


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

ubuntu 14.04 LTS iTorch 설치  (2) 2015.04.18
ubuntu 14.04 LTS Torch 설치  (0) 2015.04.17
안드로이드 가상 머신 지니모션(Genymotion) 설치  (0) 2015.04.09
멀티 스레드  (0) 2014.12.13
파이썬 argparse  (0) 2014.11.08
,

안드로이드 가상 머신 지니모션(Genymotion) 설치

프로그래밍 2015. 4. 9. 23:18

상당히 빠른 속도, 다양한 버전의 안드로이드 가상머신을 지원하는 

안드로이드 가상머신 Genymotion을 설치하는법을 소개합니다.



1. https://www.genymotion.com/ 사이트를 방문하면 아래와 같은 화면을 볼수 있습니다.


2. 화면 중앙에 있는 Get Genymotion을 클릭한뒤, 다운로드 버튼을 눌러줍니다.




3. 저는 윈도우개발 환경이므로 운영체제로 윈도우를 선택하였습니다.

   안드로이드 스튜디오, 이클립스와 연동시켜 안드로이드 앱 개발시 가상머신(virtual machine)으로 사용할 수 있습니다.




4. 다운로드를 눌러 설치파일을 받아서 실행시킨 후 다음다음다다음... 을 클릭하면 설치가 완료됩니다.








5. 가상머신을 돌리기 위해서는 Virtual Box가 필요해 또 설치해 줍니다.







6. 설치 완료!



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

ubuntu 14.04 LTS Torch 설치  (0) 2015.04.17
Lua 프로그래밍  (0) 2015.04.16
멀티 스레드  (0) 2014.12.13
파이썬 argparse  (0) 2014.11.08
[STL] List  (0) 2014.03.06
,

c++ strtok

프로그래밍/C/C++ 2014. 12. 13. 16:08


1
2
3
4
5
6
7
8
9
    char s[] = "Hello world good day to die";
    // strtok => "Hello\0world\0good\0day\0to\0die" 으로 변경됨
 
    for (char *p = strtok(s, " "); p != 0; p = strtok(0, " "))
    {
        printf("%s\n", p);
    }
 
    printf("%s\n", s);


'프로그래밍 > C/C++' 카테고리의 다른 글

스마트 포인터  (0) 2014.11.28
메모리 관리 함수(memset, memcpy, memmove)  (0) 2014.10.24
C++ 파일 입출력  (0) 2014.02.24
함수 주요 형태(_stdcall, _cdecl, _fastcall)  (0) 2014.02.24
비트 연산자  (0) 2014.02.12
,

멀티 스레드

프로그래밍 2014. 12. 13. 10:38

멀티 스레드 프로그래밍


1. 프로그램 vs 프로세스

   프로그램 : 하드 디스크에 저장되어있는 실행 코드

   프로세스 : 수행 중인 프로그램의 인스턴스


   Ctrl + F7 = Compile -> .o, .obj

   Ctrl + F5 = Compile + Linking -> .exe


1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
#if 0
 
#include <iostream>
 
using namespace std;
 
int main()
 
{
 
    cout << "Hello MultiThread..." << endl;
 
}
 
#endif 


 2. 프로세스 구성 요소

   1. 커널 오브젝트, PCB(Process Control Block) - 구조체

       - 구조체 타입 :   리눅스 = task_struct , 윈도우 = EPROCESS


    2. 코드와 데이터를 수용하는 주소 공간(메모리)

       : 프로세스는 동립된 자기만의 주소공간을 가지고 있다.


      프로세스는 자력으로 수행될 수 없다. 프로세스가 무엇인가를 수행하기 위해서는

      프로세스의 컨텍스트 내에서 수행되는 스레드가 존재해야 한다.


      OS는 프로세스를 만들때 한개의 스레드를 만들어 준다.

      : 메인 스레드, primary thread, initial thread


    3. 스레드

        : 프로세스 내에서의 제어 흐름(하나의 실행 단위)


    4. 스레드의 구성 요소

   1) 커널 오브젝트, TCB(Thread Control Block) 

           리눅스 : task_struct, 윈도우 : ETHREAD

       2) 자신만의 스텍을 가지고 있다.


    5. 멀티 스레드 모델과 멀티 프로세스 모델

        - 멀티 프로세스 모델 

        장점 : 다른 프로세스에서 잘못된 메모리 접근을 하더라도, 다른 프로세스는 영향을 받지 않음

        단점 : 데이터의 공유가 어렵다. 운영체제에서 제공하는 IPC 메커니즘을 이용해야 한다.


       - 멀티 스레드 모델

          장점 : 같은 프로세스의 컨텍스트 안에서 여러개의 스레드가 동작하는 것이기 때문에 데이터의 공유가 쉽다.

            프로세스에 비해 다루는 정보가 적기 때문에, 컨텍스트 스위칭의 비용이 작다.

          단점 : 다른 스레드에서 잘못된 메모리 접근을 하는 순간, 프로세스가 종료되기 때문에, 안정성 확보가 어렵다.

                 데이터의 공유로 인한 경쟁 상태가 발생하기 쉽고, 동기화 매커니즘이 필요하다.


       병행성(Concurrency)  vs  병렬성(Parallelism)

        - 병행성 : 사전적 의미는 어떤 일들이 동시에 발생하는 것이다.

                       어떤 일들이 동시에 수행되는 것처럼 보이지만, 실제로는 순차적으로 동작하는 것을 의미한다.

                       병행성은 단일 프로세서 시스템에서 프로세스나 스레드가 동작하는 것

        - 병렬성 : 사전적 의미는 어떤 일들이 동시에 진행되는 병행적인 작업이다.

                       작업들이 동일한 방향으로 교차됨 없이 독립적으로 진행되는 것을 의미한다.

                       병렬성은 멀티 프로세서 시스템에서 프로세스나 스레드가 동시에 동작하는 것


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

Lua 프로그래밍  (0) 2015.04.16
안드로이드 가상 머신 지니모션(Genymotion) 설치  (0) 2015.04.09
파이썬 argparse  (0) 2014.11.08
[STL] List  (0) 2014.03.06
[STL] Vector  (0) 2014.03.05
,

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
,

스마트 포인터

프로그래밍/C/C++ 2014. 11. 28. 03:31

동적 메모리 할당 과 해제에서 개발자의 실수로 메모리 누수(memory leak) 등이 발생하는 확률이 높다.

프로그램 규모가 커질수록 클래스의 내부 멤버들이 많아지고, 이를 동적 메모리로 할당해서 관리하면,

new와 delete를 자주 사용하게 된다. new로 할당받은 객체속 메모리는 반드시 delete로 해제하여야 한다.

그렇지 않으면 메모리 누수가 발생하게 된다.


C++11 표준에서 소개하는 스마트 포인터를 사용하면 편리하게 메모리 누수를 관리할 수 있다.


예제1. 스마트포인터 auto_ptr 사용

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
#include <iostream>
#include <memory>
#include <string>
using namespace std;
 
class food
{
    string name; // 이름
    int price; // 가격
 
public:
  // 생성자
    food(string _name, int _price){        
        this->name = _name;
        this->price = _price;
        cout << "생성자 호출" << endl;
    }
  // 소멸자
    ~food() {
        cout << "소멸자 호출" << endl;
    }
  
  // 정보 출력
    void printData() {
        cout << "name = " << this->name << ", price = " << this->price << endl;
    }
};
 
int main()
{
    auto_ptr<food> test(new food("apple", 30));
    test->printData();
    return 0;
}


결과


스마트 포인터 auto_ptr를 이용한 동적메모리 할당에서는 개발자가 따로 delete를 하지 않아도 

소멸자를 호출하면서 할당된 메모리를 자동으로 해제한다.


만약 동일 객체를 두 개의 포인터로 가리키게 하면 두 개의 포인터 각각의 메모리를 해제하였을 때,

중복으로 메모리를 해제하는 경우가 발생해 오류가 발생하게 된다.


메모리를 자동으로 해제해 주는 auto_ptr를 사용하면 편리하지만 한편으로는 개발자에게 매우 위엄한 상황이

발생되기도 한다. 엄밀히 말하면 C++11 표준에서는 auto_ptr 사용을 권고하지 않는다. 그렇다면 해결책이 없는것일까?


이 경우 unique_ptr를 사용하면 된다.


unique_ptr는 C++11 표준에서 사용할 수 있는 스마트 포인터이다.

unique_ptr는 auto_ptr와 유사하게 가리키는 객체의 소유권을 그대로 강조하지만, 메모리 중복 해제로

인한 런타임 오류가 발생하지는 않는다. 컴파일러가 미리 이런 오류를 처리해 준다.


unique_ptr를 사용하여 어느 정도 스마트 포인터의 요구 사항을 충족하긴 했지만, 

하나의 객체를 여러 스마트 포인터로 가리키길 원할 수 있다.

이 경우 컴파일 타임이나, 런타임에서 오류 없는 스마트 포인터를 C++11에서 제공한다.

shared_ptr라는 스마트 포인터를 사용하면 된다.




'프로그래밍 > C/C++' 카테고리의 다른 글

c++ strtok  (0) 2014.12.13
메모리 관리 함수(memset, memcpy, memmove)  (0) 2014.10.24
C++ 파일 입출력  (0) 2014.02.24
함수 주요 형태(_stdcall, _cdecl, _fastcall)  (0) 2014.02.24
비트 연산자  (0) 2014.02.12
,
1 2 3 4 5

티스토리툴바