Programming/openCL & CUDA2012. 6. 5. 07:38
Nvidia cuda 공식 문서를 찾아보니, 메모리별 소요 클럭에 대한 내용이 자세히 나온다.

local memory 나 device memory 와 같이 off-chip memory(그래픽 카드에 내장되지 않은)는 400~800 클럭에 접근하고
shared memory나 register와 같이 on-ship memory(그래픽 카드에 내장된)는 10~20 클럭 정도에 접근이 가능하다.
constant memory는 2.0 에서 부터는 1회에 대해서 400~800 클럭이고, caching이 되면서 10~20 클럭에 읽어올 수 있다.

다르게 말하면 매우 빈번하게(클럭 빈도로는 40회 이상?) 읽어오는 데이터의 경우에는
무조건 shared memory에 읽어와서 빠르게 읽는게 전체 실행 속도에 유리할 것으로 보인다.
(읽어와서 다시 저장을 해야 한다면 80회 이상으로 늘어날지도?)

5.2  Maximize Utilization 
5.2.3  Multiprocessor Level 

If all input operands are registers, latency is caused by register dependencies, i.e. some of the input operands are written by some previous instruction(s) whose execution has not completed yet. In the case of a back-to-back register dependency  (i.e. some input operand is written by the previous instruction), the latency is equal to the execution time of the previous instruction and the warp schedulers must schedule instructions for different warps during that time. Execution time varies  depending on the instruction, but it is typically about 22 clock cycles for devices of compute capability 1.x and 2.x and about 11 clock cycles for devices of compute capability 3.0, which translates to 6 warps for devices of compute capability 1.x and  22 warps for devices of compute capability 2.x and higher (still assuming that warps execute instructions with maximum throughput, otherwise fewer warps are needed).  For devices of compute capability 2.1 and higher, this is also assuming enough instruction-level parallelism so that schedulers are always able to issue pairs of instructions for each warp. 

 If some input operand resides in off-chip memory, the latency is much higher: 400 to 800 clock cycles. The number of warps required to keep the warp schedulers busy during such high latency periods depends on the kernel code and its degree of  instruction-level parallelism. In general, more warps are required if the ratio of the number of instructions with no off-chip memory operands (i.e. arithmetic instructions most of the time) to the number of instructions with off-chip memory  operands is low (this ratio is commonly called the arithmetic intensity of the program). If this ratio is 15, for example, then to hide latencies of about 600 clock cycles, about 10 warps are required for devices of compute capability 1.x and about  40 for devices of compute capability 2.x and higher (with the same assumptions as in the previous paragraph).

5.3
Maximize Memory Throughput
The first step in maximizing overall memory throughput for the application is to minimize data transfers with low bandwidth.
That means minimizing data transfers between the host and the device, as detailed in Section 5.3.1, since these have much lower bandwidth than data transfers between global memory and the device.
That also means minimizing data transfers between global memory and the device by maximizing use of on-chip memory: shared memory and caches (i.e. L1/L2 caches available on devices of compute capability 2.x and higher, texture cache and constant cache available on all devices).



[출처 : CUDA_C_Programming_Guide.pdf]

Posted by 구차니
Programming/openCL & CUDA2012. 6. 3. 22:07
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
-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
리플 룩 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]   


Posted by 구차니
Programming/openCL & CUDA2012. 6. 2. 17:00
아놔... 또 신버전이야 ㅠ.ㅠ

 

'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
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 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++ 템플렛 관련 정리가 잘되어 있는 링크 발견!

[링크 : http://ikpil.com/category/IT책%20정리실/C++%20Template# ]
Posted by 구차니
Programming/processing2012. 5. 4. 19:14
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 구차니