Programming/openCL & CUDA2012. 9. 20. 21:58
CUDA 책을 읽다가 결합법칙에서 계속 막혔었는데..
생각을 해보니 "공유 메모리"는 일정 갯수의 쓰레드 끼리만 공용으로 사용한다라는 사실을 잊고 있었던것...
다르게 보면, 공유 메모리로 복사할때는 4byte(int)형으로만 복사하면
상위레벨에서 블럭으로 system memory에서 블럭단위로 전송하여
쓰레드 블럭에서 알아서 분배하는 스타일로 복사하는 것이다.

그런 이유로, 공유 메모리 예제에서는
for문으로 왕창 복사하는게 아니라 __shared__로 정의된 배열중 하나의 값만 복사를 해서 넣는것 -_-



다르게 말하면 성능 저하를 감수하고
공유 메모리를 사용하지 않는다면 굳이 결합법칙에 머리 아플 이유도 없다는게 되려나? 

---
2012.9.22
다시보니 공유 메모리로가 아니라
로컬 메모리에서 로딩하는 모든 연산에 대한 문제이다.
cuda의 특성상 로컬 메모리(오프칩/저속) 에서 읽어오때 블럭단위로 전송을 하기에
단순하게 로컬 메모리에서 읽어 로컬메모리에 쓸때에도
로컬 메모리에서 읽는 부분의 성능 저하를 최소화 하기 위해
결합법칙을 지켜주는 것이 좋다.
 

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

cuda deviceQuery on GTX650  (0) 2013.02.17
cuda 5.0  (0) 2013.02.16
cudaMalloc 시작 위치?  (0) 2012.07.11
cudemMemcpy()  (0) 2012.06.07
cuda 에서 device memory의 용량을 초과하는 malloc은 위험해!  (0) 2012.06.06
Posted by 구차니
Programming/openCL & CUDA2012. 6. 6. 08:18
예제파일 따라한다고 512 * 65535 개의 쓰레드를 계산하게 하는 int형 배열을 무려 3개나 할당!
4byte * 512 * 65535 = 대략 128MB?

아무튼 이걸 3개를 할당하니 378MB ..
근데 ION에다가 256MB만 할당해 놓은 시스템에서 저걸 돌리니
정상처럼 돌아가는데 결과는 전부 쓰레기값(전부 결과가 0이 나옴 -_-)
그리고 미친척(!) 7번 정도 실행하니 X윈도우까지 맛이 가서 ssh로 재시작 시키게 하는 센스 OTL



결론 : 메모리 사용량은 확실히 계산하고 malloc 해주자! 

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

cudaMalloc 시작 위치?  (0) 2012.07.11
cudemMemcpy()  (0) 2012.06.07
nvcc 는 int main()을 좋아해  (0) 2012.06.05
cuda 메모리별 접근 소요 클럭 사이클  (0) 2012.06.05
ubuntu 에서 vectorAdd 직접 컴파일 하기  (0) 2012.06.03
Posted by 구차니
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. 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/openCL & CUDA2012. 5. 4. 19:07
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
일단 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
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 구차니
Programming/openCL & CUDA2012. 4. 9. 23:31
아.. 머가 먼소리인지 모르겠어서 일단 정리중..
결론 : 비싸고 새로 나온게 좋은거다~

   Multi Processor  Cores   Total
8800 GT   14  8   112
8800 GTX  16  8  128
GTX 480  15  32   480 

[링크 : http://pastebin.com/KMUXqmTY] GTX 480
[링크 : http://gpucoder.livejournal.com/990.html] 8800 GTX
2011/01/18 - [Programming/openCL / CUDA] - CUDA 3.1과 3.2의 devicequery 결과 차이점  8800 GT

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

cuda 4.2 devicequey  (0) 2012.04.23
cuda 4.2 released  (0) 2012.04.22
AMD APP SDK 예제 컴파일  (0) 2012.03.12
openCL - ATI APP SDK  (0) 2012.03.11
ATI Stream 하드웨어 요구사항  (0) 2011.10.07
Posted by 구차니
Programming/openCL & CUDA2011. 7. 3. 18:27
이넘들 끝판왕으로 맥스웰을 꺼내는건가?!


[링크 :  http://cuda.tistory.com/entry/Nvidia-케플러-맥스웰-아키텍쳐]
[링크 : http://www.maingearforums.com/showthread.php?4080-CUDA-GPU-Roadmap]

아무튼, 테그라 쪽은 슈퍼영웅(웨인=배트맨 , 로간=엑스맨, 스타크=아이언맨)을 이름으로 삼더니


[링크 : http://blogs.nvidia.com/2011/02/tegra-roadmap-revealed-next-chip-worlds-first-quadcore-mobile-processor/]
[링크 : http://chitsol.com/entry/5년 뒤에도 엔비디아는 그래픽 프로세서 기업일까?]

심심해서 KAL-EL을 검색해보니 Superboy. 그러니까 Superman의 어릴적 이름인듯?
[링크 : http://en.wikipedia.org/wiki/Superboy_(Kal-El)]

----
NVIDIA 홈페이지 간김에 가보니 2011년 5월에 4.0 버전이 릴리즈 되었다.
[링크 : http://developer.nvidia.com/cuda-toolkit-40

Latest Release

CUDA Toolkit 4.0

Archived Releases 
CUDA Toolkit 4.0 RC2 (April 2011)
CUDA Toolkit 3.2 (November 2010)
CUDA Toolkit 3.1 (June 2010)
CUDA Toolkit 3.0 (March 2010)
OpenCL 1.0 Public Release (September 2009)
CUDA Toolkit 2.3  (June 2009)
CUDA Toolkit 2.2  (May 2009)
CUDA Toolkit 2.1  (January 2009)
CUDA Toolkit 2.0  (August 2008)
CUDA Toolkit 1.1  (December 2007)

CUDA Toolkit 1.0 (June 2007)

[링크 : http://developer.nvidia.com/cuda-toolkit-archive]
 


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

openCL - ATI APP SDK  (0) 2012.03.11
ATI Stream 하드웨어 요구사항  (0) 2011.10.07
CUDA 4.0 RC  (4) 2011.03.02
CUDA - Multi GPU 프로그래밍  (1) 2011.02.05
CUDA driver API / runtime API  (0) 2011.01.26
Posted by 구차니
Programming/openCL & CUDA2011. 3. 2. 18:23
3.2도 못해봤는데 4.0이라니 언제따라가?!?!?! ㅠ.ㅠ

Release Highlights

Easier Application Porting

  • Share GPUs across multiple threads
  • Use all GPUs in the system concurrently from a single host thread
  • No-copy pinning of system memory, a faster alternative to cudaMallocHost()
  • C++ new/delete and support for virtual functions
  • Support for inline PTX assembly
  • Thrust library of templated performance primitives such as sort, reduce, etc.
  • NVIDIA Performance Primitives (NPP) library for image/video processing
  • Layered Textures for working with same size/format textures at larger sizes and higher performance

Faster Multi-GPU Programming

  • Unified Virtual Addressing
  • GPUDirect v2.0 support for Peer-to-Peer Communication

New & Improved Developer Tools

  • Automated Performance Analysis in Visual Profiler
  • C++ debugging in cuda-gdb
  • GPU binary disassembler for Fermi architecture (cuobjdump)

Please refer to the Release Notes and Getting Started Guides for more information.


[링크 : http://developer.nvidia.com/object/cuda_4_0_RC_downloads.html]


Posted by 구차니