'Programming'에 해당되는 글 1789건

  1. 2012.06.02 CUDA devicequery - ION 330
  2. 2012.06.02 cuda 5 preview
  3. 2012.05.22 nvidia ion cuda core와 h.264 library
  4. 2012.05.21 윤년 계산하기 2
  5. 2012.05.18 CUDA API 메모리 종류
  6. 2012.05.12 c++ template
  7. 2012.05.04 processing (언어)
  8. 2012.05.04 Interoperability (상호운용성)
  9. 2012.04.30 cuda 내장변수
  10. 2012.04.26 kernel block 과 thread
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 구차니
Programming/openCL & CUDA2012. 4. 30. 23:07

cuda 내장변수

일단 grid는 배재하고 생각을 하자면.
블럭과 쓰레드는 3차원으로 구성이 가능하다.
즉, 차원의 인덱스와 카운트 값이 존재해야 하는데
#define으로 하면 좋겠지만 cuda 는 그래픽 카드에서 도는 넘이다 보니
이러한 문자상수나 선언문을 넘겨줄수가 없는 구조이다.

그래서 쓰레드에게 전체 배열의 크기를 알려주기 위한 방법으로
blockDim / blockIdx / threadIdx 변수가 존재한다. 

간단하게 2차원(20x30))으로 구성된 블럭이 존재한다면
blockDim.x는 블럭의 가로 차수인 20을 모든 쓰레드에서 읽어오고
blockDim.y는 블럭의 세로 차수인 30을 모든 쓰레드에서 읽어오고 

블럭별 쓰레드에서는
blockIdx.x가 0에서 19까지
blockIdx.y가 0에서 29까지 존재하게 된다.


2011/01/16 - [Programming/openCL / CUDA] - CUDA 내장변수 - built in variable

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

CUDA API 메모리 종류  (0) 2012.05.18
Interoperability (상호운용성)  (0) 2012.05.04
kernel block 과 thread  (0) 2012.04.26
cuda 4.2 devicequey  (0) 2012.04.23
cuda 4.2 released  (0) 2012.04.22
Posted by 구차니
Programming/openCL & CUDA2012. 4. 26. 22:35

kernel block 과 thread

vectorAdd.cu 파일을 보다보니 느낌표가 똭!
    // Invoke kernel
    int N = 50000;
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
    VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N); 


kernel 에서
앞은 Grid당 Block의 갯수
뒤는 Block당 thread의 갯수를 나타낸다.

솔찍히 단일 GPU를 쓴다면 grid는 저~~~언혀 고려하지 않아도 되는데
괜히 머리 아프게 grid랑 이상한 개념들을 다 이야기 하는 바람에 이해만 어려웠던 듯 하다.


아무튼,  Devicequery를 다시 보면
블럭당 쓰레드의 최대 갯수는 512 이고
그리드당 블럭의 최대 갯수는 3차원 배열로 512x512x64가 한계이다.
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 하드웨어 구조   

그런 이유로, 아래의 예제에서는 3차원 배열로 쓰레드를 구성한 총갯수가 512를 넘지 않으면 작동을 했던 것이다.
dim3 blocksPerGrid(1,1);
dim3 threadsPerBlock(8,8,8);
이 코드는 8*8*8 = 512로 쓰레드의 최대 갯수를 넘지 않아 실행이 되지만

dim3 blocksPerGrid(1,1);
dim3 threadsPerBlock(9,9,9);
이 코드는 9*9*9 = 729로 쓰레드의 최대 갯수를 넘어 실행이 되지 않고 오류가 발생한다. 

2011/01/22 - [Programming/openCL / CUDA] - CUDA 관련 해외글   


한줄요약 : 단일 그래픽 카드로 CUDA를 하면 grid는 잊자! 좀 꺼져줘!!!!

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

Interoperability (상호운용성)  (0) 2012.05.04
cuda 내장변수  (0) 2012.04.30
cuda 4.2 devicequey  (0) 2012.04.23
cuda 4.2 released  (0) 2012.04.22
CUDA 장치별 cuda core 갯수  (0) 2012.04.09
Posted by 구차니

티스토리툴바