Programming/openCL & CUDA2012. 6. 3. 22:07

ubuntu 에서 vectorAdd 직접 컴파일 하기

SDK 에 들어있는 예제를 통채로 make 해도 되지만 하나만 직접 해보려고 하니 은근 귀찮다 -_-

/usr/local/cuda/bin/nvcc                                                << cuda 컴파일러
~/NVIDIA_GPU_Computing_SDK/C/common/inc  << include path
~/NVIDIA_GPU_Computing_SDK/shared/inc           << include path

nvcc로 컴파일 하면 아래와 같이 두번의 에러가 발생하는데 직접 파일을 검색하니 위의 include path 두개가 걸려 나왔다.
vectorAdd.cu:24:82: fatal error: sdkHelper.h: 그런 파일이나 디렉터리가 없습니다
vectorAdd.cu:25:23: fatal error: shrQATest.h: 그런 파일이나 디렉터리가 없습니다 

nvcc 야 PATH로 환경변수 잡아주면 간단해지지만
저넘의 INCLUDE_PATH가 은근 뒷통수 칠 듯..
$ /usr/local/cuda/bin/nvcc vectorAdd.cu \
   -I/home/minimonk/NVIDIA_GPU_Computing_SDK/C/common/inc \
   -I/home/minimonk/NVIDIA_GPU_Computing_SDK/shared/inc

헤더 파일은 아래의 두 곳에
~/NVIDIA_GPU_Computing_SDK/C/common/inc$ ll
합계 572
drwxrwxr-x 4 minimonk minimonk  4096  6월  2 19:00 ./
drwxrwxr-x 6 minimonk minimonk  4096  6월  2 20:25 ../
drwxrwxr-x 2 minimonk minimonk  4096  6월  2 19:00 GL/
-rw-rw-r-- 1 minimonk minimonk 14391  6월  2 19:00 bank_checker.h
-rw-rw-r-- 1 minimonk minimonk 15841  6월  2 19:00 cmd_arg_reader.h
-rw-rw-r-- 1 minimonk minimonk  3999  6월  2 19:00 cudaGLHelper.h
-rw-rw-r-- 1 minimonk minimonk  6442  6월  2 19:00 cudaHelper.h
-rw-rw-r-- 1 minimonk minimonk 20787  6월  2 19:00 cuda_drvapi_dynlink.c
-rw-rw-r-- 1 minimonk minimonk 46935  6월  2 19:00 cutil.h
-rw-rw-r-- 1 minimonk minimonk  3183  6월  2 19:00 cutil_gl_error.h
-rw-rw-r-- 1 minimonk minimonk  3492  6월  2 19:00 cutil_gl_inline.h
-rw-rw-r-- 1 minimonk minimonk  1116  6월  2 19:00 cutil_inline.h
-rw-rw-r-- 1 minimonk minimonk   953  6월  2 19:00 cutil_inline_bankchecker.h
-rw-rw-r-- 1 minimonk minimonk 13007  6월  2 19:00 cutil_inline_drvapi.h
-rw-rw-r-- 1 minimonk minimonk 17779  6월  2 19:00 cutil_inline_runtime.h
-rw-rw-r-- 1 minimonk minimonk 36730  6월  2 19:00 cutil_math.h
-rw-rw-r-- 1 minimonk minimonk 10864  6월  2 19:00 drvapi_error_string.h
drwxrwxr-x 2 minimonk minimonk  4096  6월  2 19:00 dynlink/
-rw-rw-r-- 1 minimonk minimonk 15321  6월  2 19:00 dynlink_d3d10.h
-rw-rw-r-- 1 minimonk minimonk  6202  6월  2 19:00 dynlink_d3d11.h
-rw-rw-r-- 1 minimonk minimonk  1809  6월  2 19:00 error_checker.h
-rw-rw-r-- 1 minimonk minimonk  5035  6월  2 19:00 exception.h
-rw-rw-r-- 1 minimonk minimonk 26819  6월  2 19:00 helper_cuda.h
-rw-rw-r-- 1 minimonk minimonk  5973  6월  2 19:00 helper_cuda_drvapi.h
-rw-rw-r-- 1 minimonk minimonk  4151  6월  2 19:00 helper_cuda_gl.h
-rw-rw-r-- 1 minimonk minimonk  1132  6월  2 19:00 helper_functions.h
-rw-rw-r-- 1 minimonk minimonk 22291  6월  2 19:00 helper_image.h
-rw-rw-r-- 1 minimonk minimonk 10491  6월  2 19:00 helper_string.h
-rw-rw-r-- 1 minimonk minimonk 15757  6월  2 19:00 helper_timer.h
-rw-rw-r-- 1 minimonk minimonk  1323  6월  2 19:00 multithreading.h
-rw-rw-r-- 1 minimonk minimonk  7228  6월  2 19:00 nvGLWidgets.h
-rw-rw-r-- 1 minimonk minimonk  4646  6월  2 19:00 nvGlutWidgets.h
-rw-rw-r-- 1 minimonk minimonk  2967  6월  2 19:00 nvMath.h
-rw-rw-r-- 1 minimonk minimonk 10850  6월  2 19:00 nvMatrix.h
-rw-rw-r-- 1 minimonk minimonk 12347  6월  2 19:00 nvQuaternion.h
-rw-rw-r-- 1 minimonk minimonk  6415  6월  2 19:00 nvShaderUtils.h
-rw-rw-r-- 1 minimonk minimonk 20642  6월  2 19:00 nvVector.h
-rw-rw-r-- 1 minimonk minimonk 15917  6월  2 19:00 nvWidgets.h
-rw-rw-r-- 1 minimonk minimonk  5309  6월  2 19:00 param.h
-rw-rw-r-- 1 minimonk minimonk  2321  6월  2 19:00 paramgl.h
-rw-rw-r-- 1 minimonk minimonk  1024  6월  2 19:00 rendercheck_d3d10.h
-rw-rw-r-- 1 minimonk minimonk  1027  6월  2 19:00 rendercheck_d3d11.h
-rw-rw-r-- 1 minimonk minimonk   998  6월  2 19:00 rendercheck_d3d9.h
-rw-rw-r-- 1 minimonk minimonk  7662  6월  2 19:00 rendercheck_gl.h
-rw-rw-r-- 1 minimonk minimonk 26682  6월  2 19:00 sdkHelper.h
-rw-rw-r-- 1 minimonk minimonk  1116  6월  2 19:00 stopwatch.h
-rw-rw-r-- 1 minimonk minimonk  1897  6월  2 19:00 stopwatch_base.h
-rw-rw-r-- 1 minimonk minimonk  2804  6월  2 19:00 stopwatch_base.inl
-rw-rw-r-- 1 minimonk minimonk 11608  6월  2 19:00 stopwatch_functions.h
-rw-rw-r-- 1 minimonk minimonk  4647  6월  2 19:00 stopwatch_linux.h
-rw-rw-r-- 1 minimonk minimonk  9342  6월  2 19:00 string_helper.h

~/NVIDIA_GPU_Computing_SDK/shared/inc$ ll
합계 192
drwxrwxr-x 4 minimonk minimonk  4096  6월  2 19:00 ./
drwxrwxr-x 6 minimonk minimonk  4096  6월  2 20:25 ../
drwxrwxr-x 2 minimonk minimonk  4096  6월  2 19:00 GL/
-rw-rw-r-- 1 minimonk minimonk 15439  6월  2 19:00 cmd_arg_reader.h
drwxrwxr-x 2 minimonk minimonk  4096  6월  2 19:00 dynlink/
-rw-rw-r-- 1 minimonk minimonk  5035  6월  2 19:00 exception.h
-rw-rw-r-- 1 minimonk minimonk  1323  6월  2 19:00 multithreading.h
-rw-rw-r-- 1 minimonk minimonk  7228  6월  2 19:00 nvGLWidgets.h
-rw-rw-r-- 1 minimonk minimonk  4646  6월  2 19:00 nvGlutWidgets.h
-rw-rw-r-- 1 minimonk minimonk  2966  6월  2 19:00 nvMath.h
-rw-rw-r-- 1 minimonk minimonk 10850  6월  2 19:00 nvMatrix.h
-rw-rw-r-- 1 minimonk minimonk 12347  6월  2 19:00 nvQuaternion.h
-rw-rw-r-- 1 minimonk minimonk  6108  6월  2 19:00 nvShaderUtils.h
-rw-rw-r-- 1 minimonk minimonk 20642  6월  2 19:00 nvVector.h
-rw-rw-r-- 1 minimonk minimonk 15917  6월  2 19:00 nvWidgets.h
-rw-rw-r-- 1 minimonk minimonk  8092  6월  2 19:00 rendercheckGL.h
-rw-rw-r-- 1 minimonk minimonk  6801  6월  2 19:00 shrQATest.h
-rw-rw-r-- 1 minimonk minimonk 33202  6월  2 19:00 shrUtils.h 

so 파일은 아래의 경로에 존재한다.
/usr/local/cuda/lib$ ll
합계 394240
drwxr-xr-x  2 root root      4096  6월  2 18:47 ./
drwxr-xr-x 12 root root      4096  6월  2 18:47 ../
lrwxrwxrwx  1 root root        14  6월  2 18:46 libcublas.so -> libcublas.so.4*
lrwxrwxrwx  1 root root        18  6월  2 18:47 libcublas.so.4 -> libcublas.so.4.2.9*
-rwxr-xr-x  1 root root 105568932  6월  2 18:46 libcublas.so.4.2.9*
lrwxrwxrwx  1 root root        14  6월  2 18:46 libcudart.so -> libcudart.so.4*
lrwxrwxrwx  1 root root        18  6월  2 18:47 libcudart.so.4 -> libcudart.so.4.2.9*
-rwxr-xr-x  1 root root    427344  6월  2 18:46 libcudart.so.4.2.9*
lrwxrwxrwx  1 root root        13  6월  2 18:47 libcufft.so -> libcufft.so.4*
lrwxrwxrwx  1 root root        17  6월  2 18:46 libcufft.so.4 -> libcufft.so.4.2.9*
-rwxr-xr-x  1 root root  29333272  6월  2 18:47 libcufft.so.4.2.9*
lrwxrwxrwx  1 root root        13  6월  2 18:46 libcuinj.so -> libcuinj.so.4*
lrwxrwxrwx  1 root root        17  6월  2 18:46 libcuinj.so.4 -> libcuinj.so.4.2.9*
-rwxr-xr-x  1 root root    157120  6월  2 18:46 libcuinj.so.4.2.9*
lrwxrwxrwx  1 root root        14  6월  2 18:46 libcurand.so -> libcurand.so.4*
lrwxrwxrwx  1 root root        18  6월  2 18:46 libcurand.so.4 -> libcurand.so.4.2.9*
-rwxr-xr-x  1 root root  27434820  6월  2 18:47 libcurand.so.4.2.9*
lrwxrwxrwx  1 root root        16  6월  2 18:46 libcusparse.so -> libcusparse.so.4*
lrwxrwxrwx  1 root root        20  6월  2 18:46 libcusparse.so.4 -> libcusparse.so.4.2.9*
-rwxr-xr-x  1 root root 188374460  6월  2 18:47 libcusparse.so.4.2.9*
lrwxrwxrwx  1 root root        11  6월  2 18:46 libnpp.so -> libnpp.so.4*
lrwxrwxrwx  1 root root        15  6월  2 18:46 libnpp.so.4 -> libnpp.so.4.2.9*
-rwxr-xr-x  1 root root  52381048  6월  2 18:46 libnpp.so.4.2.9* 

'Programming > openCL & CUDA' 카테고리의 다른 글

nvcc 는 int main()을 좋아해  (0) 2012.06.05
cuda 메모리별 접근 소요 클럭 사이클  (0) 2012.06.05
CUDA devicequery - ION 330  (0) 2012.06.02
cuda 5 preview  (0) 2012.06.02
nvidia ion cuda core와 h.264 library  (0) 2012.05.22
Posted by 구차니
Programming/openGL2012. 6. 2. 23:38

openGL gcc에서 컴파일 하기

-l(소문자 L/ Library) 에 GL GLU glut를 넣으면
gcc를 통해 openGL 컴파일시 라이브러리 링크를 해준다

$ gcc -lglut -lGLU -lGL 


오랫만에 다시 컴파일 하려니 다 까먹고 이게 머야 ㅠ.ㅠ
2011/09/07 - [Linux/Ubuntu] - ubuntu 에서 openGL 프로그래밍하기

'Programming > openGL' 카테고리의 다른 글

openGL state variables  (0) 2013.12.12
openGL에서 AVI 동영상 재생하기  (0) 2013.04.09
glsl과 glew의 연관관계  (0) 2011.11.22
GLSL 관련 링크  (0) 2011.11.20
GLSL 함수목록 정리  (0) 2011.11.20
Posted by 구차니
Programming/openCL & CUDA2012. 6. 2. 22:00

CUDA devicequery - ION 330

리플 룩 ion330 모델에 내장된 ion에 대한 devicequery이다.
2개의 MP가 존재해서 총 16개의 CUDA core가 존재한다.

~/NVIDIA_GPU_Computing_SDK/C/bin/linux/release$ ./deviceQuery
[deviceQuery] starting...

./deviceQuery Starting...

 CUDA Device Query (Runtime API) version (CUDART static linking)

Found 1 CUDA Capable device(s)

Device 0: "ION"
  CUDA Driver Version / Runtime Version          4.2 / 4.2
  CUDA Capability Major/Minor version number:    1.1
  Total amount of global memory:                 254 MBytes (266010624 bytes)
  ( 2) Multiprocessors x (  8) CUDA Cores/MP:    16 CUDA Cores
  GPU Clock rate:                                1100 MHz (1.10 GHz)
  Memory Clock rate:                             800 Mhz
  Memory Bus Width:                              64-bit
  Max Texture Dimension Size (x,y,z)             1D=(8192), 2D=(65536,32768), 3D=(2048,2048,2048)
  Max Layered Texture Size (dim) x layers        1D=(8192) x 512, 2D=(8192,8192) x 512
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       16384 bytes
  Total number of registers available per block: 8192
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  768
  Maximum number of threads per block:           512
  Maximum sizes of each dimension of a block:    512 x 512 x 64
  Maximum sizes of each dimension of a grid:     65535 x 65535 x 1
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             256 bytes
  Concurrent copy and execution:                 No with 0 copy engine(s)
  Run time limit on kernels:                     No
  Integrated GPU sharing Host Memory:            Yes
  Support host page-locked memory mapping:       Yes
  Concurrent kernel execution:                   No
  Alignment requirement for Surfaces:            Yes
  Device has ECC support enabled:                No
  Device is using TCC driver mode:               No
  Device supports Unified Addressing (UVA):      No
  Device PCI Bus ID / PCI location ID:           3 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 4.2, CUDA Runtime Version = 4.2, NumDevs = 1, Device = ION
[deviceQuery] test results...
PASSED

> exiting in 3 seconds:
3...2...1...done! 

그나저나.. 대역폭에서 내장형 그래픽이라 메인메모리를 공유하는데 왜 대역폭에서 이렇게 차이가 날까?
~/NVIDIA_GPU_Computing_SDK/C/bin/linux/release$ ./bandwidthTest
[bandwidthTest] starting...

./bandwidthTest Starting...

Running on...

 Device 0: ION
 Quick Mode

 Host to Device Bandwidth, 1 Device(s), Paged memory
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     887.0

 Device to Host Bandwidth, 1 Device(s), Paged memory
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     735.9

 Device to Device Bandwidth, 1 Device(s)
   Transfer Size (Bytes)        Bandwidth(MB/s)
   33554432                     5345.2

[bandwidthTest] test results...
PASSED

> exiting in 3 seconds: 3...2...1...done! 

8800GT 에 비하면 확실히 nbody 에서의 연산속도와 fps가 많이 떨어지는 느낌
(8800GT에서는 150fps에 50GFLOP/s 정도 나옴)

2010/11/02 - [Programming/openCL / CUDA] - CUDA 예제파일 실행결과 + SLI


+ 리눅스에서 nvidia 드라이버 버전 보는 방법
$ cat /proc/driver/nvidia/version
NVRM version: NVIDIA UNIX x86 Kernel Module  295.40  Thu Apr  5 21:28:09 PDT 2012
GCC version:  gcc version 4.6.3 (Ubuntu/Linaro 4.6.3-1ubuntu5)

[링크 : http://www.nvnews.net/vbulletin/showthread.php?t=127289]   


'Programming > openCL & CUDA' 카테고리의 다른 글

cuda 메모리별 접근 소요 클럭 사이클  (0) 2012.06.05
ubuntu 에서 vectorAdd 직접 컴파일 하기  (0) 2012.06.03
cuda 5 preview  (0) 2012.06.02
nvidia ion cuda core와 h.264 library  (0) 2012.05.22
CUDA API 메모리 종류  (0) 2012.05.18
Posted by 구차니
Programming/openCL & CUDA2012. 6. 2. 17:00

cuda 5 preview

아놔... 또 신버전이야 ㅠ.ㅠ

 

'Programming > openCL & CUDA' 카테고리의 다른 글

ubuntu 에서 vectorAdd 직접 컴파일 하기  (0) 2012.06.03
CUDA devicequery - ION 330  (0) 2012.06.02
nvidia ion cuda core와 h.264 library  (0) 2012.05.22
CUDA API 메모리 종류  (0) 2012.05.18
Interoperability (상호운용성)  (0) 2012.05.04
Posted by 구차니
Programming/openCL & CUDA2012. 5. 22. 20:59

nvidia ion cuda core와 h.264 library

2개의 mp를 내장하고 있는데
음. H.264 인코딩/디코딩 라이브러리 제한이 몇개부터드라...

scalar processor는 cuda core인데 1개의 mp에 8개 있고, 4개의 mp가 32 scalar 프로세서이니 
h.264 인코딩은 ion에서는 불가능 할 것으로 보인다

MPEG-2/VC-1 support 
 Decode Acceleration for G8x, G9x (Requires Compute 1.1 or higher) 
 Full Bitstream Decode for MCP79, MCP89, G98, GT2xx, GF1xx 
MPEG-2 CUDA accelerated decode with a GPUs with 8+ SMs (64 CUDA cores).  (Windows) 
 Supports HD (1080i/p) playback including Bluray content 
 R185+ (Windows), R260+ (Linux) 
 
H.264/AVCHD support 
 Baseline, Main, and High Profile, up to Level 4.1  
 Full Bitstream Decoding in hardware including HD (1080i/p) Bluray content 
 Supports B-Frames, bitrates up to 45 mbps 
 Available on NVIDIA GPUs:  G8x, G9x, MCP79, MCP89, G98, GT2xx, GF1xx 
 R185+ (Windows), R260+ (Linux) 

[출처 : CUDA_VideoDecoder_Library.pdf] 

 Supported on all CUDA-enabled GPUs with 32 scalar processor cores or more 
[출처 : CUDA_VideoEncoder_Library.pdf] 
[링크 :  http://www.vpac.org/files/GPU-Slides/05.CudaOptimization.pdf ]

Device 0: "ION"
  CUDA Driver Version:                                             2.30
  CUDA Runtime Version:                                           2.30
  CUDA Capability Major revision number:                 1
  CUDA Capability Minor revision number:                 1
  Total amount of global memory:                                 268435456 bytes
  Number of multiprocessors:                                     2
  Number of cores:                                                         16
  Total amount of constant memory:                         65536 bytes
  Total amount of shared memory per block:         16384 bytes
  Total number of registers available per block: 8192
  Warp size:                                                                     32
  Maximum number of threads per block:             512
  Maximum sizes of each dimension of a block:   512 x 512 x 64
  Maximum sizes of each dimension of a grid:     65535 x 65535 x 1
  Maximum memory pitch:                                           262144 bytes
  Texture alignment:                                                     256 bytes
  Clock rate:                                                                   1.10 GHz
  Concurrent copy and execution:                                 No
  Run time limit on kernels:                                     No
  Integrated:                                                                   Yes
  Support host page-locked memory mapping:         Yes
  Compute mode:                                                           Default (multiple host threads can use this devi
 
[링크 :  http://forums.nvidia.com/index.php?showtopic=100288 ]  

[링크 : http://www.nvidia.com/object/picoatom_specifications.html ]
[링크 : http://en.wikipedia.org/wiki/Nvidia_Ion ]

'Programming > openCL & CUDA' 카테고리의 다른 글

CUDA devicequery - ION 330  (0) 2012.06.02
cuda 5 preview  (0) 2012.06.02
CUDA API 메모리 종류  (0) 2012.05.18
Interoperability (상호운용성)  (0) 2012.05.04
cuda 내장변수  (0) 2012.04.30
Posted by 구차니
Programming2012. 5. 21. 09:52

윤년 계산하기

년/월/일 변경시 윤년인지 확인하는 귀차니즘이 필요한데
윤년 확인하는 방법을 간단하게 구현해 놓은 곳이 있어 링크!

if(((year % 400) == 0) || ((year % 4) == 0) && ((year % 100) != 0)) 
{
    // 윤년
}
else
{
     // 평년
} 
[링크 :  http://blog.daum.net/sualchi/13719905]

그레고리력의 윤년 규칙은 정확히 4년마다 윤년이 오는 율리우스력을 수정한 것이다. 정확한 규칙은 다음과 같다.
  1. 서력 기원 연수가 4로 나누어 떨어지는 해는 우선 윤년으로 한다. (2004년, 2008년, 2012년…)
  2. 100으로 나누어 떨어지는 해는 평년으로 한다. (2100년, 2200년, 2300년…)
  3. 400으로 나누어 떨어지는 해는 다시 윤년으로 한다. (1600년, 2000년, 2400년 …)

[링크 :  http://ko.wikipedia.org/wiki/윤년] 

만약 달력 입력하는 곳이 있으면
년이 변경될 경우 2월일 경우에 날짜 확인 (time_t 구조체 상으로 tm_mon 0부터 시작하므로, 2월은 1임)
월이 변경될 경우 2월에 대해서 날짜 확인
일이 변경될 경우 매월에 대한 최대 날짜를 확인해야 한다. 

---
2015.02.04

if( month==2 && ( year%4==0 && year%100!=0 || year%400==0 ) ) {
    maxDay = 29;
}


'Programming' 카테고리의 다른 글

apple 차세대 언어 swift  (0) 2014.06.03
ARToolKit / openVRML  (0) 2012.12.25
TBB/IPP  (2) 2012.02.12
프로그래밍 언어에 대한 생각  (2) 2012.01.25
S language  (0) 2011.07.01
Posted by 구차니
Programming/openCL & CUDA2012. 5. 18. 21:22

CUDA API 메모리 종류

CUDA device에서 제공하는 메모리의 종류는 다음과 같다.

5.3.2  Device Memory Accesses .................................................................... 70 
    5.3.2.1  Global Memory ............................................................................ 70 
    5.3.2.2  Local Memory .............................................................................. 72 
    5.3.2.3  Shared Memory ........................................................................... 72 
    5.3.2.4  Constant Memory ........................................................................ 73 
     5.3.2.5  Texture and Surface Memory ........................................................ 73  

[출처 :  CUDA C Programming guide.pdf] 

Local memory 와 Global memory는 그래픽 카드의 비디오 메모리(통상 512MB / 1기가 이런식으로 말하는)에 존재하고
Shared memory는 GPU 내의 Multi-Processor에 통합되어있다.

Devicequery를 비교하면서 보자면
8800GT 512MB 짜리에서
Global memory와 Local memory는 512MB 까지 가능하며
Shared memory는 블럭당 16KB 까지 가능하다.

 Device 0: "GeForce 8800 GT"
  CUDA Driver Version:                           3.20
  CUDA Runtime Version:                          3.10
  CUDA Capability Major revision number:         1
  CUDA Capability Minor revision number:         1
  Total amount of global memory:                 536543232 bytes
  Number of multiprocessors:                     14
  Number of cores:                               112
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       16384 bytes
  Total number of registers available per block: 8192
  Warp size:                                     32
  Maximum number of threads per block:           512
  Maximum sizes of each dimension of a block:    512 x 512 x 64
  Maximum sizes of each dimension of a grid:     65535 x 65535 x 1   

2011/01/02 - [Programming/openCL / CUDA] - deviceQuery on 8600GT 512MB + CUDA 하드웨어 구조      
 

예제로 들어있는 행렬곱 예제에서
shared memory를 사용하고 사용하지 않는 차이점은 아래의 그림처럼
Global memory에 직접 한 바이트씩 읽어서 계산하는지

아니면 global memory의 블럭을
shared memory로 일정 영역만(블럭 사이즈 만큼) 복사해서 계산을 하는지의 차이점이 있다.

다른 책에 의하면 global memory는 700~900 cuda clock에 읽어오고
shared memory는 거의 1 cuda clock에 읽어 온다고 하니
되도록이면 shared memory에 복사해서 더욱 빠르게 연산하는게 유리하다고 한다.

 

  
// Matrices are stored in row-major order: 
// M(row, col) = *(M.elements + row * M.width + col) 
typedef struct { 
    int width; 
    int height; 
    float* elements; 
} Matrix; 
 
// Thread block size 
#define BLOCK_SIZE 16 
 
// Forward declaration of the matrix multiplication kernel 
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix); 
 


























 
// Matrix multiplication - Host code 
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE 
void MatMul(const Matrix A, const Matrix B, Matrix C) 
{ 
    // Load A and B to device memory 
    Matrix d_A; 
    d_A.width = A.width; d_A.height = A.height; 
    size_t size = A.width * A.height * sizeof(float); 
    cudaMalloc(&d_A.elements, size); 
    cudaMemcpy(d_A.elements, A.elements, size, 
               cudaMemcpyHostToDevice); 
    Matrix d_B; 
    d_B.width = B.width; d_B.height = B.height; 
    size = B.width * B.height * sizeof(float); 
    cudaMalloc(&d_B.elements, size); 
    cudaMemcpy(d_B.elements, B.elements, size, 
               cudaMemcpyHostToDevice); 
 
    // Allocate C in device memory 
    Matrix d_C; 
    d_C.width = C.width; d_C.height = C.height; 
    size = C.width * C.height * sizeof(float); 
    cudaMalloc(&d_C.elements, size); 
 
    // Invoke kernel 
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); 
    dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y); 
    MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C); 
 
    // Read C from device memory 
    cudaMemcpy(C.elements, Cd.elements, size, 
               cudaMemcpyDeviceToHost); 
 
    // Free device memory 
    cudaFree(d_A.elements); 
    cudaFree(d_B.elements); 
    cudaFree(d_C.elements); 
} 

// Matrix multiplication kernel called by MatMul() 
__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) 
{ 
    // Each thread computes one element of C 
    // by accumulating results into Cvalue 
    float Cvalue = 0; 
    int row = blockIdx.y * blockDim.y + threadIdx.y; 
    int col = blockIdx.x * blockDim.x + threadIdx.x; 
    for (int e = 0; e < A.width; ++e) 
        Cvalue += A.elements[row * A.width + e] 
                * B.elements[e * B.width + col]; 
    C.elements[row * C.width + col] = Cvalue; 
} 












































 
 // Matrices are stored in row-major order: 
// M(row, col) = *(M.elements + row * M.stride + col) 
typedef struct { 
    int width; 
    int height; 
    int stride;  
    float* elements; 
} Matrix; 
 
// Get a matrix element 
__device__ float GetElement(const Matrix A, int row, int col) 
{ 
    return A.elements[row * A.stride + col]; 
} 
 
// Set a matrix element 
__device__ void SetElement(Matrix A, int row, int col, 
                           float value) 
{ 
    A.elements[row * A.stride + col] = value; 
} 
 
// Get the BLOCK_SIZExBLOCK_SIZE sub-matrix Asub of A that is 
// located col sub-matrices to the right and row sub-matrices down 
// from the upper-left corner of A 
__device__ Matrix GetSubMatrix(Matrix A, int row, int col)  
{ 
    Matrix Asub; 
    Asub.width    = BLOCK_SIZE; 
    Asub.height   = BLOCK_SIZE; 
    Asub.stride   = A.stride; 
    Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row 
                                         + BLOCK_SIZE * col]; 
    return Asub; 
} 
 
// Thread block size 
#define BLOCK_SIZE 16 
 
// Forward declaration of the matrix multiplication kernel 
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix); 
 
// Matrix multiplication - Host code 
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE 
void MatMul(const Matrix A, const Matrix B, Matrix C) 
{ 
    // Load A and B to device memory 
    Matrix d_A; 
    d_A.width = d_A.stride = A.width; d_A.height = A.height; 
    size_t size = A.width * A.height * sizeof(float); 
    cudaMalloc(&d_A.elements, size); 
    cudaMemcpy(d_A.elements, A.elements, size, 
               cudaMemcpyHostToDevice); 
    Matrix d_B; 
    d_B.width = d_B.stride = B.width; d_B.height = B.height; 
    size = B.width * B.height * sizeof(float); 
    cudaMalloc(&d_B.elements, size); 
    cudaMemcpy(d_B.elements, B.elements, size, 
               cudaMemcpyHostToDevice); 
 
    // Allocate C in device memory 
    Matrix d_C; 
    d_C.width = d_C.stride = C.width; d_C.height = C.height; 
    size = C.width * C.height * sizeof(float); 
    cudaMalloc(&d_C.elements, size); 
 
    // Invoke kernel 
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE); 
    dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y); 
    MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C); 
 
    // Read C from device memory 
    cudaMemcpy(C.elements, d_C.elements, size, 
               cudaMemcpyDeviceToHost); 
 
    // Free device memory 
    cudaFree(d_A.elements); 
    cudaFree(d_B.elements); 
    cudaFree(d_C.elements); 
} 
 
// Matrix multiplication kernel called by MatMul() 
__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C) 
{ 
    // Block row and column 
    int blockRow = blockIdx.y; 
    int blockCol = blockIdx.x; 
 
    // Each thread block computes one sub-matrix Csub of C 
    Matrix Csub = GetSubMatrix(C, blockRow, blockCol); 
     // Each thread computes one element of Csub 
    // by accumulating results into Cvalue 
    float Cvalue = 0; 
 
    // Thread row and column within Csub 
    int row = threadIdx.y; 
    int col = threadIdx.x; 
 
    // Loop over all the sub-matrices of A and B that are 
    // required to compute Csub 
    // Multiply each pair of sub-matrices together 
    // and accumulate the results 
    for (int m = 0; m < (A.width / BLOCK_SIZE); ++m) { 
 
        // Get sub-matrix Asub of A 
        Matrix Asub = GetSubMatrix(A, blockRow, m); 
 
        // Get sub-matrix Bsub of B 
        Matrix Bsub = GetSubMatrix(B, m, blockCol); 
 
        // Shared memory used to store Asub and Bsub respectively 
        __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; 
        __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; 
 
        // Load Asub and Bsub from device memory to shared memory 
        // Each thread loads one element of each sub-matrix 
        As[row][col] = GetElement(Asub, row, col); 
        Bs[row][col] = GetElement(Bsub, row, col); 
 
        // Synchronize to make sure the sub-matrices are loaded 
        // before starting the computation 
        __syncthreads(); 
 
        // Multiply Asub and Bsub together 
        for (int e = 0; e < BLOCK_SIZE; ++e) 
            Cvalue += As[row][e] * Bs[e][col]; 
 
        // Synchronize to make sure that the preceding 
        // computation is done before loading two new 
        // sub-matrices of A and B in the next iteration 
        __syncthreads(); 
    } 
 
    // Write Csub to device memory 
    // Each thread writes one element 
    SetElement(Csub, row, col, Cvalue); 
}


'Programming > openCL & CUDA' 카테고리의 다른 글

cuda 5 preview  (0) 2012.06.02
nvidia ion cuda core와 h.264 library  (0) 2012.05.22
Interoperability (상호운용성)  (0) 2012.05.04
cuda 내장변수  (0) 2012.04.30
kernel block 과 thread  (0) 2012.04.26
Posted by 구차니
Programming/C++ STL2012. 5. 12. 10:45

c++ template

C++ 템플렛 관련 정리가 잘되어 있는 링크 발견!

[링크 : http://ikpil.com/category/IT책%20정리실/C++%20Template# ]

'Programming > C++ STL' 카테고리의 다른 글

unsigned char -> int 변환 주의사항  (0) 2013.02.04
템플릿 메타프로그래밍  (0) 2013.01.06
리눅스에서 c++ 컴파일시 strcpy / strcat 오류  (0) 2011.10.13
new / new[] / delete / delete[]  (4) 2010.09.16
cout 그리고 namespace  (0) 2010.09.16
Posted by 구차니
Programming/processing2012. 5. 4. 19:14

processing (언어)

10년 정도 된 언어라는데 가시화를 위한 언어이자 IDE라고 한다.
위키피디아를 보니 아두이노용 IDE랑 같은것 같은데 흐음...

[링크 : http://en.wikipedia.org/wiki/Processing_(programming_language)]
[링크 : http://processing.org/ ] 

'Programming > processing' 카테고리의 다른 글

processing 실행  (0) 2014.01.12
프로세싱(언어)  (0) 2013.01.08
Posted by 구차니
Programming/openCL & CUDA2012. 5. 4. 19:07

Interoperability (상호운용성)

CUDA를 보다보면
openGL / DirectX와의 interoperability 라는 용어가 나온다.
딱히 번역하기에는 애매해서 단어를 찾아보니 상호운용(interoperate)라는 말이 나온다.
굳이 붙여도 문제는 없어 보이지만..
[링크 : http://endic.naver.com/enkrEntry.nhn?entryId=62756ad640c44b41919ec9e5b504d898]  

3.2.11  Graphics Interoperability 
Some resources from OpenGL and Direct3D may be mapped into the address space of CUDA, either to enable CUDA to read data written by OpenGL or Direct3D, or to enable CUDA to write data for consumption by OpenGL or Direct3D. 

3.2.11 그래픽 상호운용성
OpenGL이나 Direct3D에 의해 쓰여진 데이터를 CUDA에서 읽도록 허용 하거나
OpenGL 이나 Direct3D에 의해서 사용(소비)될 데이터를 CUDA가 쓸 수 있도록 
OpenGL과 Direct3D로 부터의 몇몇 자원들은 CUDA의 주소 공간에 할당(연관/매핑)되어질수 있다. 

의역하자면, "CUDA와 OpenGL / Direct3D와의 상호연계를 위해 메모리 통로를 양쪽으로 연결해 준다." 정도이려나?

'Programming > openCL & CUDA' 카테고리의 다른 글

nvidia ion cuda core와 h.264 library  (0) 2012.05.22
CUDA API 메모리 종류  (0) 2012.05.18
cuda 내장변수  (0) 2012.04.30
kernel block 과 thread  (0) 2012.04.26
cuda 4.2 devicequey  (0) 2012.04.23
Posted by 구차니

티스토리툴바