Vietnamen’s Weblog

Time, Chances, Diligence, Intelligence: which is the most important?

Archive for the ‘Tin Sinh học’ Category

Phần 6: Fortran + CUDA (2) variables

leave a comment »

Variables

Như ta đã đề cập, CUDA có nhiều cấp độ bộ nhớ, vì thế, việc chọn lựa nơi lưu trữ cũng rất quan trọng tới hiệu năng của chương trình.

  1. register (thanh ghi): (read/write) per-thread
  2. local memory: (read/write) per-thread
  3. shared memory (read/write) per block
  4. global memory (read/write) per grid
  5. constant memory (read-only) per grid
  6. texture memory (read-only) per grid

Mặc định, khi khai báo biến ở host subprogram/modules, dữ liệu được lưu trữ  ở host global memory.

Tối đa 1 trong các thuộc tính sau được khai báo với biến: device, constant, shared, pinned

Device

REAL, DEVICE :: dev_var, x(100)

real :: y
attributes(device) :: y

Một device variable có dữ liệu được lưu ở device global memory. Nó có thể được khai báo ở host subprogram hoặc device subprogram.

Nếu khai báo trong module, nó có thể được truy cập bởi mọi device subprogram trong module đó, và mọi host subproram trong module đó hoặc là subprogram có sử dụng module

Nếu khai báo trong 1 host subprogram, thì nó có thể được truy xuất từ chính subprgram đó, hoặc các subprogram chứa trong (contains) nó.

Device variables (arrays) không thể

  • có các thuộc tính POINTER, TARGET, ALLOCATABLE.
  • xuất hiện trong COMMON block
  • xuất hiện trong EQUIVALENCE statement
  • là thành viên của một kiểu mới (derived data type)

Biến hằng có thể làm đối số (actual arguments) cho một host subprogram hoặc device subprogram, với điều kiện INTERFACE tường minh, và đối số lúc định nghĩa (dummy arguments) tương ứng cũng có thuộc tính DEVICE.

Nếu là kiểu mảng (array) nó có thể là explicit-shape array, allocatable device variable, hoặc (nếu ở trong host subprogram) assumed-shaped dummy array.

Constant

Có 2 cách khai báo

real :: c(100)
attributes(constant) :: c

hoặc

real, constant :: d(100)

Dữ liệu hằng không thể

  • có các thuộc tính POINTER, TARGET, ALLOCATABLE.
  • xuất hiện trong COMMON block
  • EQUIVALENCE statement
  • là thành viên của một kiểu mới (derived data type)

Mảng (array) là hằng thì phải có kích thước cố định.

Biến hằng có thể làm đối số (actual arguments) cho một host subprogram hoặc device subprogram, với điều kiện INTERFACE tường minh, và đối số lúc định nghĩa (dummy arguments) tương ứng cũng có thuộc tính CONSTANT.

Dữ liệu hằng:

  • không thể thay đổi nếu định ra trong device subprogram
  • có thể thay đổi nội dung, nếu định ra trong host subprogram

Shared

Biến chia sẻ chỉ có thể lưu dữ liệu trên shared memory space. Vì thế, nó chỉ được khai báo trong device subprogram, và cũng chỉ có thể được truy xuất từ chính device subprogram đó, hoặc các device subprogram khác mà bản thân biến này được truyền tới dưới dạng đối số (actual argument)

có thể bị đọc/ghi từ mọi threads trong block. CHÚ Ý: khi một thread cập nhật dữ liệu lên 1 shared variable, để đảm bảo là khi thread khác đọc sẽ đọc dữ liệu mới này, ta cần đồng bộ sau khi cập nhật dữ liệu, dùng SYNCTHREADS() intrinsic function.

Dữ liệu chia sẻ không thể

  • có các thuộc tính POINTER, TARGET, ALLOCATABLE.
  • xuất hiện trong COMMON block
  • EQUIVALENCE statement
  • là thành viên của một kiểu mới (derived data type)

Mảng (array) là SHARED thì có thể dùng asumed-sized array nếu nó không  phải được khai báo là 1 dummy argument của 1 device subprogram. Vậy kích thước là bao nhiêu? điều này sẽ được qui định khi gọi kernel dùng chevron syntax với đối số thứ 3 (bytes) chỉ định số bytes ở shared memory dành cho assumed-sized array. Nếu có nhiều hơn một assumed-size array làm đối số, chúng sẽ được ngầm định là tương đương nhau, i.e. cùng bắt đầu ở cùng địa chỉ ô nhớ. Vì thế, nếu dùng nhiều hơn 1 assumed-sized array, nhà lập trình cần quản lí nó.

attributes(global) subroutine sub(y)
   real, shared :: x(4, *)
   integer :: y 
   ...
end subroutine

Nếu mảng (array) là SHARED và không phải là dummy argument + không phải assumed-sized array, thì nó phải là fixed-sized array.

Biến chia sẻ có thể làm đối số (actual arguments) cho một host subprogram hoặc device subprogram, với điều kiện INTERFACE tường minh, và đối số lúc định nghĩa (dummy arguments) tương ứng cũng có thuộc tính SHARED.

Pinned

Chỉ dành để khai báo với biến thuộc kiểu allocatable array và khai báo trong module/host subprogram.

real, allocatable, dimension(:), pinned :: x

Khi cấp phát, dữ liệu được lưu ở host page-locked memory. Ưu điểm là dữ liệu lưu ở đây sẽ được copy tới device memory nhanh hơn so với từ host normal-paged memory tới device memory. Vì giới hạn về bộ nhớ dành cho host page-locked memory là tùy hệ thống, dữ liệu không đảm bảo lưu trữ ở đây nếu vượt quá kích thước cho phép, lúc đó nó sẽ tự động lưu vào host normal-paged memory.

Pinned array có thể dùng làm đối số cho host subprogram dù INTERFACE của nó có được khai báo hay không, cũng như không phụ thuộc vào dummy argument có được khai báo là PINNED hay ALLOCATABLE hay không. Tuy nhiên, nếu giải phóng nó (deallocate) bên trong subprogram, thì dummy argument phải có khai báo là PINNED, nếu không sẽ có lỗi.

Tham khảo:

  1. CUDA Fortran programming guide and reference

Written by vietnamen

Tháng Mười 8, 2009 at 11:21 sáng

Phần 5: Fortran + CUDA (1) subprogram

leave a comment »

subprogram (Subroutine/Function)

Define

Chương trình chính sẽ luôn chạy trên host, còn các subprograms của nó có thể chạy hoặc trên host hoặc trên device. Để một subprogram của Fortran có thể chạy trên device, nó cần sử dụng thêm các từ chỉ định attributes(device)attributes(global) đi trước các từ khóa subroutine, function. Khi đó, subprogram còn có tên gọi là  1 device kernel.

  1. Subroutine: có thể dùng attributes(host), attributes(global), attributes(device) hoặc cả attributes(host, device)
  2. Function: có thể dùng host, device, hoặc cả hostdevice

NOTE:

  • Function không thể có global attribute.
  • Subprogram với attributes(hostdevice) phải nằm trong 1 module
! Host subroutine

subroutine solve( n, a, x, y )
   real, device, dimension(*) :: x, y
   real :: a
   integer :: n

! call the kernel
   call ksaxpy<<<n/64, 64>>>( n, a, x, y )
end subroutine

Invoke

Để gọi 1 kernel từ host function ta dùng cú pháp đặc biệt <<<num_block, num_thread [, bytes [, streamid]]>>> gọi là chevron syntax. Trong ví dụ trên, n chính là số lượng threads cần chạy song song cho kernel này và cũng chính là kích thước của vectors xy. Ở đây, số threads trong mỗi block luôn là 64 nên số lượng blocks là n/64.

CHÚ Ý: Nếu kích thước lớn hơn số threads song song tối đa, ta sẽ phải có chỉnh sửa đoạn code.

! Kernel definition
attributes(global) subroutine ksaxpy( n, a, x, y )
real, dimension(*) :: x,y
real, value :: a
integer, value :: n, i
i = (blockidx%x-1) * blockdim%x + threadidx%x
if( i <= n ) y(i) = a * x(i) + y(i)
end subroutin
! Kernel definition
! Scalar Alpha X Plus Y: y = ax + y
attributes(global) subroutine ksaxpy( n, a, x, y )
   real, dimension(*) :: x,y
   real, value :: a
   integer, value :: n, i

   i = (blockidx%x-1) * blockdim%x + threadidx%x
   if( i <= n ) y(i) = a * x(i) + y(i)
end subroutine

Passing dummy arguments

Với device subprograms, các dummy arrguments (scalar, array…) được MẶC ĐỊNH là truyền tham khảo (passed by reference), i.e. truyền địa chỉ – không truyền nội dung, nếu biến bị thay đổi bên trong subprogram thì giá trị bên ngoài cũng tự động được cập nhật + nơi lưu trữ là ở host memory.

Nếu đối số là giá trị đơn (scalar argument), thì có thể truyền giá trị bằng cách thêm thuộc tính VALUE vào trước khai báo biến đó (value dummy arrguments).

attributes(global) subroutine add(a, b, n)
   real, dimension(n) :: a, b
   integer, value :: n
...
end subroutine

Limitation

Device kernel (attributes(device))

  1. không được là hàm đệ qui (recursive)
  2. không chứa biến với  SAVE attribute, hay khởi tạo giá trị ban đầu (data initialization)
  3. không có đối số tùy chọn (optional arguments), i.e. đối số với giá trị khởi tạo mặc định
  4. không được là một phần (contains) trong host subprogram, subroutine hay function nào cả.
  5. bản thân nó cũng không được có chứa (contains) bất kì subroutines, functions nào.
  6. đối số (dummy argument) không được dùng assumed-shaped array và không được có pointer attribute.
  7. chỉ có thể được gọi từ host function, và dùng chevron syntax
  8. có thể có cả attributes(device) và attributes(host), lúc đó subprogram sẽ được biên dịch 2 lần, một lần cho host và một lần cho device

Kernel kernel (attributes(global))

  1. kích thước toàn bộ các đối số (dummy argument) bị giới hạn kích thước 256 bytes (nghĩa là tối đa 64 phần tử với kiểu integer(kind=4) và tối đa 32 phần tử (với kiểu integer(kind=8))
  2. không được có cả attributes(host) và attribute(global) đồng thời

Đồng thời, khi gọi subroutine đó, cần cho biết có bao nhiêu cá thể (instances) của kernel đó chạy song song. Mỗi cá thể sẽ chạy trên 1 threads độc lập. Như ta đã giới thiệu các CUDA threads được tổ chức thành từng nhóm, gọi là thread blocks. Như vậy mỗi thread sẽ có 1 global thread ID (đại diện cho thread block) và 1 local thread ID (đại diện cho thread bên trong thread block đó).

Tương tự với CUDA C, ở đây, để biết global thread ID, ta dùng biến blockidx, và local thread ID thông qua biến threadidx.

Written by vietnamen

Tháng Mười 7, 2009 at 11:37 sáng

Phần 4: Fortran + CUDA + Accelerator (0) install + compile

leave a comment »

NVIDIA phối hợp cùng PGI (Portland Group) cuối cùng đã cho ra Fortran compiler với CUDA support. Trước đây, CUDA support chỉ có trên C/C++ programming language. Hiện nay hầu hết các code chương trình cho khoa học là viết trên Forrtran, vì thế, nhu cầu chuyển dần sang CUDA-enabled Tesla NVIDIA GPU là rất lớn. Bên cạnh đó, PGI cung cấp mô hình Accelerator Programming Model (APM) nhằm giúp chuyển đổi các source codes có sẵn sang GPU với ít thay đổi nhất.

Đăng kí bản quyền

Để dùng được PGI Fortran, bạn cần có file license.dat, ở đây tôi không đi vào chi tiết cài đặt. Tuy nhiên, đây là bước đầu tiên và chỉ làm 1 lần.

 

Kích hoạt CUDA + accelerator compiler

Kể từ phiên bản 9.0-2, PGI Fortran hỗ trợ CUDA features. Để kích hoạt, bạn cần tạo ra 1 file có tên sitenvrc để khai báo các thư mục cài đặt CUDA software. File nỳ nằm ở thư mục $PGI/linux86-64/9.0-4/bin (CHÚ Ý: $PGI là alias đến thư mục cài đặt PGI, ví dụ: $PGI=/opt/pgi, hay $PGI=/usr/pgi, tuỳ theo version của PGI Fortran mà 9.0-1 có thể phải đổi thành 9.0-2 hay khác).

Sau đó, thêm các dòng sau vào trong file sitenvrc (chú ý dấu chấm phẩy sau mỗi hàng là quan trọng)

set NVOPEN64DIR=/opt/cuda/open64/lib;   
set CUDADIR=/opt/cuda/bin;
set CUDALIB=/opt/cuda/lib64;   ! thay lib64 bằng lib nếu là CUDA 2.2
set GCCVERSION=40301; ! có một số không tương thích giữa G++4.3 với NVIDIA installation files

Nếu dòng thứ 4 không có, bạn có thể gặp hàng loạt  lỗi như sau:

/usr/include/c++/4.3/x86_64-linux-gnu/bits/c++config.h(233): error: expected a “{“

….

Để kiểm tra cài đặt thành công, với CUDA, bạn có thể test với deviceQuery hay bandwidthTest (nằm trong bộ CUDA SDK samples), hoặc pgaccelinfo tool. Với pgaccelinfo, thông tin gồm

Device Number:                 2
Device Name:                   Tesla C1060
Device Revision Number:        1.3
Global Memory Size:            4294705152
Number of Multiprocessors:     30
Number of Cores:               240
Concurrent Copy and Execution: Yes
Total Constant Memory:         65536
Total Shared Memory per Block: 16384
Registers per Block:           16384
Warp Size:                     32
Maximum Threads per Block:     512
Maximum Block Dimensions:      512 x 512 x 64
Maximum Grid Dimensions:       65535 x 65535 x 1
Maximum Memory Pitch:          262144B
Texture Alignment              256B
Clock Rate:                    1296 MHz
Initialization time:           3376 microseconds
Current free memory            4246142976
Upload time (4MB)              1091 microseconds ( 908 ms pinned)
Download time                  1545 microseconds (1365 ms pinned)
Upload bandwidth               3844 MB/sec (4619 MB/sec pinned)
Download bandwidth             2714 MB/sec (3072 MB/sec pinned)

Chú ý, “Device Revision Number” là 1.3 và đây cũng là output mặc định khi biên dịch, muốn compile ra CUDA code theo phiên bản khác, ví dụ 1.0, ta thêm vào một dòng khác (1.1 hoặc 1.0)

set COMPUTECAP=1.0

Biên dịch với pgf95

Để biên dịch ra code chạy được trên PGPGU

  • using GPU Programming model CUDA, ta du’ng compiler flag “-Mcuda” (không có dấu ngoặc kép)
  • using Accelerator Programming model, ta du’ng compiler flag “-ta=nvidia” (không có dấu ngoặc kép)

NOTE: Để biên dịch các thư viện kèm theo với accelerator, ta cũng cần dùng compiler flag “-ta”.

Debug

Another way to edit the file sitenvrc

set NVDIR=/usr/local/cuda;
set NVOPEN64DIR=$NVDIR/open64/lib;
set CUDADIR=$NVDIR/bin;
set CUDALIB=$NVDIR/lib64;
set GCCVERSION=40303;

Tham khảo:

  1. http://www.nvidia.com/object/cuda_home.html
  2. http://www.topcoder.com/tc?module=Static&d1=sponsors&d2=NVIDIAOverview
  3. http://www.bv2.co.uk/?p=910

Written by vietnamen

Tháng Mười 4, 2009 at 10:11 chiều

Phần 3: CUDA extension convention + syntax + compiler

leave a comment »

http://stackoverflow.com/questions/242894/cuda-driver-api-vs-cuda-runtimeTại thời điềm này, CUDA đã được cấu hình xong. Ta có thể bắt đầu tìm hiểu cách biên dịch và cách viết 1 chương trình có dùng CUDA.

CUDA_architecture

Trình biên dịch

Quá trình biên dịch diễn ra qua nhiều giai đoạn:

  1. code cho CPU và code cho GPU sẽ được tách ra
  2. code cho GPU sẽ được convert sang một ngôn ngữ trung gian giống như assembly language, gọi là PTX (Parallel Thread eXecution). Tại đây, nó sẽ được kiểm tra tính tối ưu. Cuối cùng, nó sẽ được chuyển về dạng mã máy mà GPU có thể hiểu ở dạng nhị phân.

compiling_CUDA

Một chương trình CUDA lúc chạy cần có cuda.so (CUDA core library), cudart.so (CUDA runtime library). Bên cạnh đó, có các thư viện viết sẵn (CUDA library): CUFFT (Fast Fourier Transform), CUBLAS (math)

architecture_CUDA

C/C++

nvcc là một command-line tool để biên dịch ứng dụng CUDA. Với ngôn ngữ C, nvcc sẽ gọi gccNVIDIA PTX compiler.

  • nvcc <filename>.cu [-o <executable>]
    • Builds release mode
  • nvcc -g <filename>.cu
    • Builds debug mode
    • Can debug host code but not device code
  • nvcc -deviceemu <filename>.cu
    • Builds device emulation mode
    • All code runs on CPU, no debug symbols
  • nvcc -deviceemu -g <filename>.cu
    • Builds debug device emulation mode
    • All code runs on CPU, with debug symbols
    • Phù hợp cho debug chương trình

Fortran

Với Fortran, ta dùng pgf95 kèm compiler flag “-Mcuda”  (chú ý: “-ta=nvidia” là dùng cho PGI Fortran acceleration model không phải CUDA)

CHÚ Ý: có thêm compiler flag “-tp” để chỉ định loại CPU mà chương trình được biên dịch sẽ chạy (nhằm tối ưu hóa code), ví dụ: amd64, amd64e…

Ta sẽ tìm hiểu kĩ hơn đối với PGI Fortran ở các bài viết sau.

Các ngôn ngữ khác:

PyCUDA: http://documen.tician.de/pycuda/tutorial.html

 

File extension

C/C++

Các hàm kernel được lưu ở files với phần mở rộng là .cu

Fortran

.cuf cho file ở dạng free-format

.CUF như .cuf nhưng nó sẽ thông qua quá trình tiền xử lí (preprocess) trước khi được compile.

Khởi tạo bộ nhớ

Như ta đã biết có nhiều phân cấp bộ nhớ khác nhau của CUDA-capable GPU. Tuỳ vào nơi muốn cấp phát, CUDA cung cấp các hàm khác nhau

DEVICE MEMORY

C/C++

  1. cudaMalloc(void** pointer, size_t nbytes)
    cấp phát nbytes (BYTEs) với địa chỉ khởi đầu lưu ở pointer
  2. cudaMemset(void* pointer, int value, size_t count)
    khởi tạo (gán) giá trị là value cho count ô nhớ, bắt đầu tại địa chỉ lưu ở pointer
  3. cudaFree(void* pointer)
    giải phóng bộ nhớ
  4. cudaMemcpy():
    sao chép dữ liệu bên trong host memory, bên trong device memory và đồng thời giữa 2 loại  (vì thời gian sao chép là giữa 2 loại là chậm, nên hạn chế sử dụng tối đa)
int n=1024;
int nbytes=1024*sizeof(int);
int * d_a = 0; ! khai báo con trỏ

cudaMalloc((void**)&d_a, nbytes);
cudaMemset(d_a, 1, nbytes);
cudaFree(d_a)

Fortran

HOST MEMORY

C/C++

  1. pointer = malloc(int nbytes)
  2. free(in*ptr)
int* h_a=0;
int dimx=16;

int num_bytes=dimx*sizeof(int);

h_a = (int) malloc(num_bytes);

if (h_a == 0) {
   print ("Could not allocate");
  return 1;
}

Fortran

Sao chép dữ liệu

C/C++

  1. cudaMemcpy (void* dst, void* src, size_t nbytes, enum cudaMemcpyKind direction)
    đảm bảo sao chép an toàn (chỉ bắt đầu sao chép khi các lệnh gọi tới CUDA đều hoàn tất, i.e. không có kernel nào đang thực hiện; đồng thời mọi thread bị blocked cho tới khi dữ liệu sao chép xong)
  2. enum cudaMemcpyKind: nhận các giá trị có thể sau
    cudaMemcpyHostToDevice
    cudaMemcpyDeviceToHost
    cudaMemcpyDeviceToDevice

Fortran

Subprogram (Functions/Subroutine)

Vì trong 1 chương trình, có subprogram thì chạy ở host, có cái thì chạy ở device. Nên để chỉ định một cách tường minh, các từ khóa (specifier) sau được dùng.

__device__   (C/C++),  attributes(device) (Fortran)

xác định kernel (chạy trên device), chỉ được gọi từ một kernel khác, không có biến tĩnh (static variables), chỉ có thể là hàm đi kèm (inlined), không hỗ trợ đệ qui (recursion)

__global__  (C/C++), attributes(global)  (Fortran)

xác định kernel (chạy trên device), chỉ có thể được gọi từ một host function, kiểu trả về luôn là void (C/C++); khi được gọi, phải chỉ rõ cấu hình thực thi kernel (execution configuration – xem ở sau) – nghĩa là cho biết số blocks trong 1 grid, số lượng threads trong 1 block.

chúng thực thi bất đồng bộ.

__host__  (C/C++), attributes(host)  (Fortran)   – mặc định

xác định host function (chạy trên host), chỉ có thể được gọi từ 1 host function khác, đây là chỉ định mặc định nếu hàm được định nghĩa mà không dùng từ khóa nào ở trên đi kèm.

Variables

__device__and__constant__

biến chỉ nằm trong global memory (device memory), tồn tại theo chương trình (chỉ mất khi chương trình kết thúc), có thể được truy cập từ mọi thread ở mọi grid.

biến hằng (constant variables) chỉ có thể được gán giá trị từ host functions.

__shared__

nằm trong block shared memory, tồn tại theo block, chỉ có thể được truy cập từ các threads của block đó. Nếu chỉ khai báo, thì không có giá trị đầu mặc định.

<none>

là các biến tự động (automatic variables), lưu ở trong thanh ghi (registers) của mỗi thread khi có thể, nếu kích thước quá lớn, thì sẽ lưu ở global memory.

Hạn chế với kernel

Kernel là đoạn code để chạy trên GPU, nó thường được tổ chức thành các function hoặc subroutine. Tuy nhiên, vẫn có một số hạn chế nhất định đối với code chạy trên GPU so với code trên CPU.

C/C++

  1. chỉ truy xuất bộ nhớ trên GPU (không được đọc data từ CPU, vì thế trước khi thực thi kernel, từ host phải thực hiện các tác vụ khởi tạo các đối tượng dữ liệu trên GPU và sao chép dữ liệu cần thiết từ host memory ra đó.
  2. mỗi function phải có số đối số (arguments) là cố định
  3. không có biễn tĩnh (static)
  4. không có đệ qui
  5. kernel phải được khai  báo dùng 1 trong các qualifier sau
    __global__                     : kernel này được gọi từ host (CPU), trả về luôn là void
    __device__                     : kernel này được gọi từ 1 kernel khác
    __host__                        : function thực thi trên CPU (mặc định nếu không chỉ định qualifier)
    __host__ __device__
__global__ void kernel_name(int *a)
{
   int idx = blockIdx.x * blockDim.x + threadIdx.x;
   a[idx] = idx;
}

Fortran

  1. a
  2. a
  3. a
  4. không có đệ qui

Gọi 1 kernel (Execution configuration)

Cú pháp:

C/C++: function_name<<<gd, bd [, N]>>>(param1, param2, param3)

Fortran: function_name<<<gd, bd [, N [, streamid] ]>>>(param1, param2, param3)

Ở đây, function_name là tên hàm cần gọi, param_i là các đối số của nó. Quan trọng cần nắm là

  • gd: kích thước của grid (kiểu dim3, với 3 thành phần: gd.x, gd.y, gd.z cho biết số blocks ở mỗi chiều
  • bd: kích thước của mỗi blocks (kiểu dim3, với 3 thành phần: bd.x, bd.y, bd.z cho biết số threads ở mỗi chiều
  • N: (optional) cho biết số lượng bộ nhớ (tính theo bytes) cần thêm vào trong shared memory của mỗi block để dùng cho đối số là assumed-sized array

Như vậy, ta thấy là để phân chia công việc cho mỗi threads, ta cần xác định các thông tin về threads. CUDA đã có các biến đặc biệt sẵn để hỗ trợ

  1. gridDim: (kiểu dim3), cho biết kích thước của grid
  2. blockIdx: (kiểu uint3?), chứa index của block trong grid
  3. blockDim: (kiểu dim3), cho biết kích thước của blocks
  4. threadIdx: (kiểu uint3), chứa index của thread trong block

C/C++

Mỗi kernel có thể đươc son song hoá bởi 1 số lượng thread nhất định (ví dụ Tesla T10 có 240 SP, thì có thể có tối đa 30,720 threads song song). Tuy nhiên, đễ dễ quản lí, CUDA tổ chức chúng thành tứng blocks và grids. Mỗi grid sẽ tương ứng thực thi cho 1 kernel. Mỗi grid hiện nay cho phép 2D các blocks, mỗi block cho phép 3D các threads. Như vậy, khi gọi thực thi 1 kernel, 2 thông tin trên là tối thiểu cần phải có. Cú pháp để chỉ định 2 thông tin này gọi là chevron syntax. Đồng thời 1 kiểu dữ liệu mới, nhằm khai báo dữ liệu 3 chiều gọi là dim3 ra đời.

dim3 grid(16,16)
dim3 block(16,16)
kernel_name<<<grid, block, 0, 0>>> (đối số)
kernel_name<<<32, 512>>> (đối số);

Thực ra có tới 4 thông tin trong chevron syntax, nhưng 2 cái cuối là tuỳ chọn, đó là:

  1. size of shared memory: 0 by default
  2. stream ID: 0 by default

Chú ý là các phân bố/sắp xếp có thể là 1-, 2- hoặc 3-D phù hợp với cấu trúc dữ liệu cần sử lí (vector, mảng 2 chiều (hình ảnh…), mảng 3 chiều). Do đó, 1 kiểu dữ liệu mới ra đời dim3.

type(dim3)
   integer(kind=4) :: x, y, z
end type

Ví dụ: đối với vector, ví dụ có 16 thành phần, thường ta chỉ dùng block và grid 1 chiều, nếu block size là 4 thì ta  sẽ dùng 4 blocks.

int dimx=16;
int num_bytes = dimx * sizeof(int);
dim3 grid, block;
block.x = 4; ! mỗi block là 4 threads
grid.x = dimx/block.x;
int*d_a = 0;
cudaMalloc((void**)&d_a, num_bytes);

kernel_name<<<grid, block>>> (int* d_a)

CHÚ Ý: Đối với C/C++, chỉ số ban đầu là 0 (zero).

như vậy, thread 1 của block 1 sẽ truy xúât phần tử thứ 0; thread 1 của block 2 sẽ truy xuất phần tử thứ 4, … Như vậy element mà thread tương ứng cấn xử lí có thể được xác định thông qua.

int idx = blockIdx.x * blockDim.x + threadIdx.x;

Ví dụ: đối với ma trận 2 chiều, ta thường chia làm các submatrix (ở đây đơn giản là chia chẵn, nếu chia không chẵn, ta cần có giải pháp tối ưu hơn sẽ nói ở 1 bài viết khác), mỗi submatrix tương ứng với 1 block, như vậy số block chính là dimension của ma trận cũ chia cho dimension của submatrix.

int dimx=16;
int dimy=20;
dim3 grid, block;
block.x = 4; ! submatrix of size 4x4
block.y = 4;
grid.x = dimx/block.x;
grid.y = dimy/block.y;
kernel_name<<<grid, block>>> (int*a, int dimx, int dimy)

như vậy, thread 1 của block 1 sẽ truy xúât phần tử thứ 0; thread 1 của block 2 sẽ truy xuất phần tử thứ 4, … Như vậy element mà thread tương ứng cấn xử lí có thể được xác định thông qua.

int ix = blockIdx.x * blockDim.x + threadIdx.x;
int iy = blockIdx.y * blockDim.y + threadIdx.y;
int idx = iy*dimx + ix

Ví dụ: n là số lượng threads thực thi cho kernel, nếu BLOCKSIZE là số lượng thread trong mỗi block thì số lượng block là n/BLOCKSIZE. Ta cần tính toán nhân matrix-vector: y=Ax

/// C/C++ code snippet
dim3 block(BLOCKSIZE)    /// 1-D
dim3 grid(n/BLOCKSIZE)    /// 1-D
mv_kernel<<<grid, block>>>(subA, subX, subY, n)

Trong C/C++, chỉ số bắt đầu là zero, vì thế block.x=0..BLOCKSIZE-1, block.y=0, block.z=0 trong trường hợp này.

Fortran

Truy xuất từ xa

Để cho phép truy xuất từ xa, bật quyền read-write cho mọi users với /dev/nv* devices.

 

References:

  1. http://stackoverflow.com/questions/242894/cuda-driver-api-vs-cuda-runtime

Written by vietnamen

Tháng Mười 4, 2009 at 10:07 chiều

Phần 2: Bắt đầu với CUDA-capable GPU

leave a comment »

Cài đặt

Cách cài đặt phần cứng + phần mềm, xin xem tài liệu đi kèm.

CUDA SDK sẽ mặc định cài đặt ở $(HOME)/NVIDIA_CUDA_SDK

CUDA Toolkit sẽ mặc định cài đặt ở /usr/local/cuda

Kiểm tra phần cứng

Để cập nhật thông tin về phần cứng, ta dùng lệnh

sudo update-pciids

Để biết CUDA-capable GPU nào đang được cài đặt, ta dùng lệnh

lspci | grep -i nVidia
81:00.0 3D controller: nVidia Corporation GT200 [Tesla C1060 / Tesla S1070] (rev a1)

0a:00.0 VGA compatible controller: nVidia Corporation NV45GL [Quadro FX 3400/4400] (rev a2)

81:00.0 3D controller: nVidia Corporation GT200 [Tesla C1060 / Tesla S1070] (rev a1)

Ở đây, ta thấy có 2 card đồ họa (NVIDIA Quadro), và 1 card NVIDIA Tesla.

Kiểm tra phần mềm

Để viết được ứng dụng song song dùng CUDA-capable GPU, ta cần CUDA Development tool (hiện nay là version 2.3) gồm có CUDA toolkit CUDA SDK.  CUDA toolkit chứa các header files, tools, libraries… để giúp biên dịch ứng dụng chạy trên nền CUDA-capable GPU. CUDA SDK chứa các hướng dẫn cần thiết cho việc viết các ứng dụng song song dùng công nghệ CUDA.

Về phần hệ điều hành, cả Windows lẫn Linux đều có hỗ trợ. Ở đây, người viết sử dụng Ubuntu Jaunty 9.04 (64-bit). Về phần trình biên dịch (compiler), để tương thích, gcc version 3.4.x hoặc 4.x.x. PGI Fortran compiler cũng vừa có hỗ trợ, tuy nhiên đây là phiên bản có bản quyền, và nó cần gcc 4.3.x hoặc 4.4.0. Trong các phần sau, PGI Fortran sẽ được nói tới khi dùng với CUDA.

Cuối cùng là NVIDIA driver. CUDA toolkit 2.3 yêu cầu NVIDIA driver version 190 trở đi. Để kiểm tra version của NVIDIA driver, dùng X-tool

/usr/bin/nvidia-settings

hoặc

cat /proc/driver/nvidia/version

 

hoặc dùng deviceQuery tool ở phần sau.

Thiết lập môi trường để có thể dùng CUDA

Nơi cài đặt mặc định với CUDA toolkit là ở /usr/local/cuda$(HOME)/NVIDIA_CUDA_SDK.

Vì thế, cần đảm bảo biến môi trường PATH có chứa /usr/local/cuda/bin. Biến môi trường LD_LIBRARY_PATH có chứa hoặc /usr/local/cuda/lib hoặc /usr/local/cuda/lib64, tùy vào môi trường là 32-bit hay 64-bit. Đưa các thông tin biến môi trường vào file ~/.bash_profile. Sau đó, cập nhật dùng

load ~/.bash_profile

CUDA SDK được cài đặt ở /usr/local/NVIDIA_GPU_Computing_SDK hoặc mặc định ở $(HOME)/NVDIA_GPU_Computing_SDK.


Kiểm tra cài đặt CUDA-capable GPU

Để kiểm tra, ta sẽ biên dịch gói SDK codes kèm theo NVIDIA_GPU_Computing_SDK/C/. Trong thư mục NVIDIA_GPU_Computing_SDK/C có chứa rất nhiều ví dụ viết trên ngôn ngữ C. Với những ai dùng C, có thể dùng đây làm nguồn tham khảo. Nếu biên dịch thấy lỗi, đọc phần tương thích và lỗi.

Sau khi được biên dịch, các binary files sẽ nằm ở NVIDIA_GPU_Computing_SDK/C/bin/linux/release. Sau đây là một số tiện ích có thể chạy test.

deviceQuery

deviceQuery là một trong các binary files vừa được biên dịch xong. Tiện ích này sẽ cho ta biết thông số của GPU card. CHÚ Ý: Nếu chức năng SELinux có sử dụng, ta cần disable nó trước khi chạy deviceQuery.

sudo setenforce 0

bandwidthTest

bandwidthTest là một trong các binary files vừa được biên dịch xong. Tiện ích này sẽ kiểm tra xem giao tiếp giữa CUDA-capable GPU và hệ thống có tốt hay không. Nó có nhiều tham số hữu ích, nên dùng –help để xem thêm.

 

Tương thích và lỗi

Trước khi biên dịch, cần dùng gcc 4.3. Hiện nay (Oct, 12, 2009) phiên bản gcc 4.4. có nhiều vấn đề với CUDA 2.3.

Lúc biên dịch,có thể thiếu một số header files, bạn có thể kiểm tra các packages sau:

Nếu thiếu file glut.h, cài đặt freeglut-dev. Nếu vẫn chưa được (hiếm khi xảy ra) ta có thể thử với các gói sau

libgl1-mesa-dev
mesa-common-dev
libgl1-mesa-dri-dev
libglu1-mesa-dev
libx11-dev

Nếu thiếu …., cài đặt libxi-devel, libxmu-dev.

Nếu bị lỗi


obj/release/deviceQueryDrv.cpp.o: In function `main':
deviceQueryDrv.cpp:(.text+0x134): undefined reference to `cuDriverGetVersion'
collect2: ld returned 1 exit status
make[1]: *** [../../bin/linux/release/deviceQueryDrv] Error 1
Ta kiểm tra version của NVDIA driver dùng

 

cat /proc/driver/nvidia/version

 

 

nếu là 180.xxx thì phải nâng cấp lên 190.xxx (để tương thích với CUDA 2.3) theo cách sau

  • NVIDIA driver 180.xxx (phiên bản kèm theo repos của Ubuntu 9.04) có thể làm việc với CUDA 2.2, nhưng không với CUDA 2.3
  • NVIDIA driver 190.xxx có thể làm việc với CUDA 2.3. Link download (nhớ disabled và remove các packages đã cài đặt với NVIDIA driver 18.xxx dùng synaptic).

Có thể remove 180.44 dùng

sudo NVIDIA-Linux-x86-180.44-pkg1.run –uninstall

Sau đó, xoá hoàn toàn các files có chứa từ nvidia ở trong các thư mục sau

/lib/modules/2.6…/ (there can be more than one folder)
/usr/lib64/xorg

Cuối cùng, tiến hành cài đặt mới với NVIDIA driver 190.xx (ví dụ: NVIDIA 190.18 Beta).

Tham khảo:

  1. http://developer.nvidia.com/object/opengl_3_driver.html
  2. http://www.nvidia.com/object/cuda_get.html
  3. http://www.nvidia.com/object/unix.html
  4. http://forums.opensuse.org/hardware/412100-nvidia-driver-180-44-opensuse-11-1-a.html

Written by vietnamen

Tháng Mười 4, 2009 at 6:29 chiều

Phần 1: CUDA programming model

leave a comment »

Như đã giới thiệu ở bài viết trước, Tesla là thế hệ GPGPU card của NVIDIA cung cấp CUDA – một tập các lệnh mở rộng của ngôn ngữ C viết theo chuẩn (industry-standard C) giúp cho việc viết các chương trình song song. Vậy thế nào là lập trình song song và kiến trúc phần cứng để hỗ trợ cho nó như thế nào? – Thông tin chi tiết có thể tham khảo ở trang này.

CPU architectures?

Mọi quá trình đều tồn tại cái gọi là “atom”, nghĩa là đơn vị nhỏ nhất không thể chia nhỏ. Trong lập trình máy tính, ban đầu, “atom” chính là các dòng mã máy. Với 1 chương trình được viết bởi 1 ngôn ngữ nào đó, để có thể thực thi, cần được biên dịch ra mã máy gồm các chuỗi nhị phân  (0 và 1) có chiều dài xác định; máy tính, cụ thể là bộ vi xử lí trung tâm (CPU – Central Processing Unit), sẽ tuần tự đọc từng dòng mã máy để thực hiện toàn bộ chương trình. CPU lấy kiến trúc của Von Neumann làm nền tảng.

Von_Neumann_architecture

Kiến trúc CPU của Von Neumann

Theo hình trên, ALU (Arithmetic Logic Unit) là nơi trực tiếp xử lí tác vụ được lấy từ mã máy. Một mã máy là một chuỗi bit nhị phân (0,1) với chiều dài xác định (8-bit, 32-bit, 64-bit – thế hệ càng về sau thì số bít càng cao) chia làm nhiều phần.

– phần đầu: chứa thông tin tác vụ cần thực hiện (op code), ví dụ: add, sub, or, xor, not, and…

– phần tiếp theo chứa địa chỉ ô nhớ nơi lưu trữ dữ liệu mà op code đó tác động lên

– …

Khi chạy 1 chương trình đã được biện dịch, phần mã máy sẽ nạp vào bộ nhớ – có tên gọi là bộ nhớ chương trình (program memory, nằm trong RAM). Sau đó, CPU sẽ thực hiện các bước

  1. đọc từng dòng lệnh (fetch),
  2. giãi mã nội dung dòng lệnh (decode),
  3. thực thi lệnh (execute), và
  4. ghi lại kết quả (writeback).

Nếu chia quá trình thực thi 1 lệnh ra các bước nhỏ hơn (step) như ở trên: fetch, decode, execute, writeback – “atom” bây giờ là 1 step, lệnh thứ (n) thực thi xong step 1 thì lệnh (n+1) có thể thực thi step 1, không cần đợi lệnh (n) thực thi tới hết step 4 như kiến trúc cũ. Điều này làm tăng tốc độ xử lí lên – đây chính là cơ chế đường ống (pipeline). Tuy vậy, cơ chế pipeline cũng không thể giúp cho cùng 1 step của 2 lệnh khác nhau chạy đồng thời. Đồng thời, mỗi lệnh máy cũng chỉ tác dụng lên 1 địa chỉ ô nhớ. Các CPU dạng này gọi là scalar processor, nó thuộc kiến trúc SISD (single-instruction, single-data).

Vì thời gian lấy lệnh (fetch) thường là khá lâu – do thời gian đọc từ RAM chậm, giải pháp bộ nhớ đệm (cache memory: L1, L2) với khả năng truy cập rất nhanh từ CPU đã ra đời để giúp lưu sẵn 1 loạt các lệnh trong cache memory, điều này giảm thời gian lấy lệnh (fetch) từ bộ nhớ chương trình.

Trong thực tế, nhu cầu xử lí nhiều dữ liệu cùng lúc, với cùng 1 tác vụ là rất phổ biến, ví dụ: nhân một ma trận cho 1 giá trị, như vậy mọi phần tử của ma trận đều được thực hiện cùng 1 tác vụ nhân với cùng 1 số. Vì đơn vị thực hiện tính toán trong CPU chính là ALU (arithmetic logic unit), việc tăng số lượng ALU trên mỗi CPU có thể giúp khả năng xử lí nhiều data cùng lúc. Đồng thời, các dòng lệnh máy thay vì chứa thông tin 1 địa chỉ ô nhớ riêng lẽ, sẽ chứa 1 vùng địa chỉ ô nhớ hay một tập hợp các địa chỉ ô nhớ. Các CPU dạng này gọi là vector processor (array processor), nó thuộc kiến trúc SIMD (single-instruction, multiple data).

Sau này thì có kiến trúc multicore (dual-core, quad-core) nằm trên cùng 1 dye.

dual_core

Dual core

Còn nhiều kiến trúc nữa nhưng ta không thể nói hết ở đây. Ngoài các bộ vi xử lí đa dụng như vậy, các dòng vi xử lí cho đồ họa cũng từng bước ra đời – gọi là GPU (Graphics Processing Unit). Trước khi nói đến GPU, ta giới thiệu sơ về threads. Một thread là một bản sao (hiện thực – instances) của một đoạn code mà được thực thi như 1 đơn vị – i.e. thực hiện từ đầu đến cuối mà không bị can thiệp. Để định danh một threads, hệ thống cần 1 hoặc nhiều thông tin (contextual information). Tùy vào đó, mà thời gian để chuyển từ thread này sang thread kia có thể nhanh (vài chục microseconds) hoặc chậm (vài ngàn microseconds). Qua đó, threads được phân loại là: heavy weight, middleweight, và lightweight.

  • heavy weight: để xác định 1 thread cần rất nhiều thông tin: hardware register, kernel stack, user-level stack, …  (Unix process)
  • middle weight: các thread chỉ nằm trong 1 không gian địa chỉ nên chỉ cần 1 thông tin để xác định ra thread
  • light weight: thread loại này không cần vùng ô nhớ riêng, chúng chia sẻ với các threads khác trong cùng 1 process nên việc chuyển đổi qua lại giữa các thread rất nhanh.

GPU architecture?

GPU là các bộ vi xử lí với các opcode được thiết kể chủ yếu cho các tác vụ đồ họa. GPU có mặt ở khắp nơi, trong các hệ thống nhúng, điện thoại di động, máy tính để bàn và các thiết bị trò chơi. Các lệnh cho 2D thông thường trong đồ họa như: BitBLT (kết hợp nhiều ảnh theo 1 toán tử nào đó – Raster operation (OR, AND, XOR…) – để cho ra ảnh kết quả), các lệnh để vẽ các hình cơ bản – primitives (chữ nhật, hình tam giác, hình tròn, hình cung). Các GPU sau này có hỗ trợ thêm cho các lệnh 3D.

Các lệnh mà các GPU cung cấp là ngôn ngữ cấp thấp (assembly language). Để tạo điều kiện cho các nhà lập trình sử dụng, các thư viện viết cho các ngôn ngữ cấp cao (C/C++…) ra đời.

  • 2D: WinG, DirectDraw, OpenGL
  • 3D: Direct3D, OpenGL

Với các thế hệ GPU về sau, nhiều lệnh phức tạp cũng được thêm vào, e.g. để thực hiện các tác vụ: H&T (hardware transform & lightning), pixel shader (fragment shader), và vertex shader. Thuật ngữ shader dùng để chỉ (1) một chương trình máy tính chạy trên một môi trường đặc biệt để (ví dụ: GPU), (2) tập hợp các lệnh trong GPU dùng để tính toán các hiệu ứng đổ bóng với độ chính xác cao. Vì yêu cầu cần thực hiện các tác vụ shader rất nhiều, nên các GPU thường có đơn vị xử lí shader chạy song song.

Nói về song song, có 3 cơ chế chính: MPI, OpenMP Stream Processing. Chức năng của GPU phù hợp với mô hình stream processing qua đó có 3 giai đoạn chính: gather, operate, và scatter. Nghĩa là dữ liệu (ví dụ: các pixels từ file hình ảnh) – có thể liên tục hoặc ở vị trí bất kì – sẽ được tập hợp lại (gather) thành 1 dòng (stream). Sau đó các một chuỗi các chương trình con (gọi là kernel functions) sẽ tác động tuần tự (operate) lên mỗi thành phần trong stream – thông thường mỗi kernel đều xử lí toàn bộ các thành phần trong stream (nên còn gọi là uniform streaming) và đầu ra của kernel function này sẽ là đầu vào của kernel function kế tiếp (nó hoạt động theo cơ chế pipeline) để giảm thiểu việc đọc/ghi dữ liệu trở lại bộ nhớ. Cuối cùng, dữ liệu đã xử lí được bởi mọi kernel functions sẽ được phân bố trở lại (scatter) bộ nhớ.

Streaming processing là cơ sở cho kĩ thuật graphics pipeline (rendering pipeline) cho phép chuyển đổi và biểu diễn từ hình ảnh dạng vectors ra hình ảnh dạng raster (pixels, dots) để có thể hiển thị/xuất ra ở các thiết bị đầu ra (màn hình, máy in, file ảnh lưu trong máy) với tốc độ cực kì nhanh (in realtime). Cùng với sự thay đổi của kiến trúc như đã đề cập ở trên, công nghệ VLSI cũng nâng tốc độ tính toán lên rất nhiều bằng cách tăng số lượng transistors trên một đơn vị diện tích.

Công nghệ chế tạo GPU

Công nghệ chế tạo GPU và số lượng transitor trên mỗi chip

GPGPU architectures?

Với sức mạnh của GPU ngày càng lấn át CPU, nhu cầu sử dụng GPU cho tính toán đa dụng (general purpose computation, cụ thể là cho các ứng dụng SIMD [single instruction, multiple data] hay vector-based datasets), không đơn thuần để tăng tốc cho các ứng dụng đồ họa (ASAP [application-specific application processor]), trở thành một vấn đề thiết thực. Từ đó khái niệm “general purpose GPU” ra đời.

Để so sánh giữa GPU vs. CPU, John Peddie, chủ tịch của John Peddie Research đã phát biểu: “The CPU cannot do the parallel operations that a GPU does, and a GPU cannot do the scalar operations that a CPU does. They need each other, and neither can replace the other.”

CPU-GPU-benchmarkGPGPU với hàng trăm cores có khả năng chạy hàng ngàn threads song song là một ứng viên tiềm tàng cho việc giải quyết các bài toán lớn có tính song song cao (real-time simulations, molecular simulations). Điều này đòi hỏi việc thêm một số lệnh máy cần thiết vào GPU. Đồng thời, để hỗ trợ việc sử dụng thuận lợi các lệnh cấp thấp của GPGPU, NVIDIA đã cung cấp một bộ thư viện mở rộng cho ngôn ngữ C, với tên gọi CUDA; và kiến trúc của các GPGPU card này gọi là Tesla. AMD/ATI thì có Brook+ nhưng không được bàn đến ở đây. Chú ý là miền ứng dụng của GPGPU không phải là cho mọi bài toán, có một số bài toán mà GPGPU không đem lại độ lợi (gain) đáng kể.

Tesla architecture & CUDA programming model

Kiến trúc logic

Ở phần này, có nhiều thuật ngữ quan trọng cần nắm. Trong CUDA programming model, CPU được xem là host, còn GPGPU được xem là device. Có thể có nhiều hosts và nhiều devices trên một máy tính. Cả hai đều có bộ nhớ riêng với tên gọi tương ứng là host memorydevice memory, i.e. không có bộ nhớ dùng chung. Vì thế, việc trao đổi dữ liệu cần copy qua lại và chỉ diễn ra khi không có đoạn code nào đang chạy trên GPGPU (để đảm biều điều này CUDA cung cấp 1 lệnh cho phép đồng bộ để đợi mọi lệnh trên GPU xong). MỘt số thử nghiệm cho thấy copy từ GPU->CPU nhanh hơn 1 ít so với CPU->GPU.

Với các đặc trưng của một GPU đã đề cập ở trên, GPGPU phù hợp để thực thi 1 tác vụ (i.e. subroutine/function) được gọi nhiều lần, mỗi lần áp dụng trên các dữ liệu đầu vào khác nhau. Đoạn code cho tác vụ như vậy khi thực thi trên GPGPU được gọi là 1 [device] kernel [function]. Khi có lệnh gọi kernel đó, chương trình sẽ kiểm tra (thông qua 1 cú pháp gọi kernel được qui định sẵn) xem có bao nhiêu hiện thực (instances, array of parallel threads)  chạy song song cho kernel đó. Điểm hạn chế hiện nay (cho tới thế hệ 2 của Tesla) là tại mỗi thời điểm chỉ có 1 kernel được thực thi, với số lượng hiện thực của kernel đó là rất lớn. Để dễ quản lí một số lượng lớn threads như vậy, cũng đồng thời giúp ánh xạ (mapping) dễ dàng từ dữ liệu – vector-based dataset – sang các threads, kiến trúc thread blocksgrids được sử dụng. Khi gọi 1 kernel thực thi, ta phải qui định tường minh dimension của block và của grid.

  • Để dễ quản lí, phân chia dữ liệu, các thread được gộp thành từng nhóm gọi là thread blocks theo kích thước 1-, 2- hoặc 3-D, và số lượng tối đa threads trong mỗi thread blocks là có giới hạn (<= 2^9 = 512 threads/block). Ví dụ: một 3D block thì có thể có kích thước là 8x8x8.
  • Để tăng số lượng thread song song có thể thực hiện cho 1 kernel, các thread blocks có thể gộp thành 1 grid [of thread blocks] với cách sắp xếp cũng có thể theo 1-, 2- hoặc 3-D (<= 2^32 blocks/grid). Như vậy, tối đa, ta có thể có 2^32 * 2^9 = 2^41 threads tổng cộng chạy song song cho 1 kernel.

Các thread của GPU là light-weight. Mỗi thread đều có dữ liệu của riêng nó (bao gồm thanh ghi – register – và local memory). Các threads trong cùng 1 block thì có thể chia sẽ dữ liệu qua shared memory, nhưng điều này không thể xảy ra với hai threads ở hai blocks khác nhau, ngoại trừ việc lưu trữ ở device memory (không nằm trên mỗi GPU chip – global memory). Tốc độ truy xuất đến dữ liệu trong shared memory là cực kì nhanh. Dữ liệu của thread lưu ở shared memory sẽ mất khi thread kết thúc. Nếu với dữ liệu quá lớn ( > 16KB), nó sẽ tự động được lưu trữ ở device memory chính là DDRAM (kích thước có thể đến 4GB, hoặc thậm chí 16GB với Tesla S1070). Dữ liệu lưu ở device memory sẽ tồn tại  trong suốt thời gian chương trình hoạt động (hoặc khi mà nó bị deallocate đối với dữ liệu động). Vì thế, dữ liệu lưu ở device memory có thể được truy xuất từ mọi kernel. Đồng thời, dữ liệu ở device memory của chương trình đang chạy sẽ là nơi trao đổi qua lới với host memory.

A grid with 1D array of threads

A grid with 1D array of thread blocks.

Grids_and_Blocks_of_Threads

A grid with 2D array of thread blocks, each block is a 2D array of threads.

Kiến trúc phần cứng

Thông tin trên là về mặt cấu trúc logic. Về mặt vật lí (phần cứng), Tesla GPGPU có các 32-bit cores (thread processors, streaming processor core, hay scalar processor- SP) – đơn vị xử lí nhỏ nhất. Số lượng là 128 cores với 1st gen. (generation), 240 cores với 2nd gen., và tương lai 512 cores với 3rd gen.

1st gen. chưa có bộ xử lí kép. Từ 2nd gen., Tesla có thêm 64-bit double precision processor (DP) nhưng số lượng bằng 1/8 so với SP. Một SP là đơn vị phần cứng để chạy tương ứng 1 thread,  i.e. mỗi SP chỉ xử lí 1 thread tại một thơi điểm mỗi SP hỗ trợ 128 threads. Cứ mỗi 8 SP, cùng với các đơn vị phần cứng khác (1 I-cache, 1 MT-issue, 1 C-cache, 2 SFU, 1 shared memory) được gộp thành một nhóm, tương ứng là 1 multiprocessors (streaming multiprocessor – SM). Rồi cứ mỗi 2 SM (hoặc nhiều hơn tuỳ vào kiến trúc), cùng với các đơn vị phần cứng khác (1 Texture unit, 1 Geometry controller (aka TPC controller), 1 SMC), tạo thành 1 group, tương ứng là 1 Thread Processing Clusters (TPC).

CUDA_SP

Mỗi 8 SP được gồm thành 1 SM. Tùy theo số lượng SP (32, 128, 240, ...) mà ta có số lượng SM khác nhau. Ở đây là kiến trúc cho GPU có 32 SP - 1st gen. Tesla.

Nói tóm lại, với 2nd gen. Tesla:

  1. texture unit gồm 24KB L1 cache (mỗi 2 SM có 1 texture unit) + ROP unit phục vụ cho các tác vụ bên đồ họa
  2. mỗi SM có
    tám 32-bit cores: có thể thực hiện 2 phép toán dấu chấm động (gồm 1 đơn vị để tính cọng + 1 đơn vị để tính nhân) trong 1 chu kì
    một 64-bit double-precision processsor (DP) có thể tính cộng/nhân với độ chính xác kép
    hai SFU cho phép tính các phép toán đặc biệt: nghịch đảo (1/x), 1/sqrt(x), sine, cosine, …
    8KB constant cache (lưu trữ các giá trị hằng số)
    instruction cache (lưu các lệnh của kernel)
    16KB shared memory + 64KB register memory
  3. mỗi special function unit (SFU) có thể thực hiện 4 phép toán dấu chấm động (4 đơn vị tính nhân) trong 1 chu kì
  4. nếu có 240 SP, và 128 threads chạy song song trong mỗi SP, thì ta có thể có tối đa 30,720 threads
  5. tỉ lệ số lượng DP trên SP là 1/8.

Về lí thuyết, nếu mọi units đều được sử dụng, mỗi SM có thể tính toán 8*2 + 2*4 = 24 tác vụ trong 1 chu kì. Hình vẽ sau cho thấy việc ánh xạ giữa phần cứng và phần mềm.

Mô hình hiện thực

Mô hình hiện thực

Khi gọi 1 kernel, nhà lập trình chỉ quan tâm là cần bao nhiêu thread blocks, và bao nhiêu thread trong mỗi block dùng để chạy kernel đó, còn việc ánh xạ thread blocks tới multiprocessor nào là việc của phần cứng, nhà lập trình không cần quan tâm. Đặc trưng này gọi là scalable programming model.

Ví dụ: NVIDIA GeForce 8 series có 128 SPs, được chia làm 16 SM. 16 SM lại được chia làm 8 đơn vị xử lí độc lập, gọi là Texture Processor Clusters (TPC). Mỗi multiprocessor có 8 single-precision streaming processors, 2 special-function units (SFU), 1 đơn vị lấy lệnh chạy đa luồng (multithreaded instruction fetch  and issue unit (MT issue)), 1 bộ nhớ đệm lưu các lệnh nạp sẵn (instruction cache), 1 bộ nhớ lưu trữ các biến hằng (read-only constant cache),  và 1 thanh ghi chia sẻ với dữ liệu tối đa 16-KB (read/write shared memory). GeForce 8 series

TPC_GeForce8

GeForce 8800 và một TPC gồm 2 SM

Tesla T8/G80 này gồm có 8 TPC. Mỗi TPC gồm có một texture unit và 2 SMs

TPC_Cuda

Ví dụ: NVIDIA GeForce 200 series có 240 SP, được chia làm 30 SM. Mỗi SM có 8 single-precision SP, 1 double-precision unit, và 1 shared memory. Như vậy, nếu dùng độ chính xác kép, mỗi SM tại mỗi thời điểm chỉ chạy 1 thread, thay vì là 8 đối với độ chính xác đơn. Do đó, tốc độ sẽ tăng 8 lần nếu dùng độ chính xác đơn. Các vấn đề khi sử dụng độ chính xác đơn hoặc kép sẽ được nói rõ ở các phần tiếp.

G80_Hardware

Tesla C1060 có 240 SP (freq: 1.3GHz, RAM:4GB, bus: PCI Expressx16 Gen.2), gồm 10 TPC (mỗi TPC có 24 SP, được gộp vào 3 SM, mỗi SM có 8 SP)

CUDA_Tesla_T10

Tesla T10 (2nd gen. Tesla)

Tesla S1070: gồm 4 Tesla T10 kết nối với nhau

Phân cấp bộ nhớ

Phân cấp bộ nhớ

Vì tốc độ chậm trong truy xuất bộ nhớ là nguyên nhân chủ yếu gây ra hiện tưỡng nghẽn cổ chai (bottleneck) trong các chương trình song song. Tesla tạo ra nhiều lớp (layer) bộ nhớ khác nhau với tốc độ truy cập khác nhau phù hợp cho các nhu cầu cụ thể. Quyền truy cập đến mỗi loại bộ nhớ cũng tuỳ vào cấp độ phần cứng sử dụng.

Các loại bộ nhớ cần biết:

  1. register (thanh ghi): (read/write) per-thread –> Mỗi SM có 8,192 registers.
  2. local memory: (read/write) per-thread  –>
  3. shared memory (read/write) per block
  4. global memory (read/write) per grid
  5. constant memory (read-only) per grid –> Mỗi SM có 8KB cho cả constant và texture
  6. texture memory (read-only) per grid

CUDA cung cấp các hàm cho phép cấp phát bộ nhớ động trên từng loại bộ nhớ cụ thể (ta sẽ tìm hiểu ở phần khác).memory_CUDA

Cơ chế phân công công việc

Cơ chế mà mỗi SM thực hiện là SIMT (single instruction multiple-threads), tương tự với cơ chế SIMD của CPU nhưng phức tạp hơn. Chi tiết cụ thể về kĩ thuật, tuy nhiên, không được công bố.

Trước hết, ta bàn về thuật ngữ thread trong CUDA. Đó đơn giản là phần tử cơ bản của dữ liệu cần xử lí, không có liên quan đến CPU threads hay GPU threads thuần túy. Thuật ngữ “warp” dùng để chỉ một nhóm 32 threads như vậy. 32 là kích thước tối thiểu để xử lí song song bởi 1 SM theo kiểu SIMD. Nói bên lề, thuật ngữ “warp” xuất phát từ “weaving”.

Để thuận lợi cho nhà lập trình, thay vì tương tác trưc tiếp với các warps, CUDA đưa ra khái niệm block, mỗi block chứa từ 64 đến 512 threads. Mỗi SM cho phép tối đa 8 blocks, tùy vào kích thước mỗi blocks miễn sao tổng số threads phải nhỏ hơn 768 (sẽ nói ở sau vì sao là 768). Vì thế, để tối ưu hóa công việc, nên phân chia số threads trong mỗi blocks là chia hết cho 32 (bạn có thể chia thành 1×32, hoặc 16×16, hoặc 8x8x16 … – kích thước mỗi block). Tiếp đến, việc nhóm các blocks vào trong 1 grid cho phép gọi và thực thi hàng ngàn threads chỉ bằng 1 lệnh gọi đơn, và không cần quan tâm đến tài nguyên của phần cứng. Nghĩa là nếu có ít SM, thì các blocks sẽ được thực thi tuần tự, còn nếu có nhiều SM thì chúng sẽ được thực thi song song. Đây chính là yếu tố làm nên đặc tính “extensible” của CUDA proramming model.

Hiện nay, 2nd gen. Tesla có SM gồm 8 SP và 1 DP. Mỗi SM có thể quản lí cùng lúc tới 24 warps, do đó cho phép tối đa 24×32=768 threads song song trên 1 SM. Như vậy, số threads trên toàn bộ GPGPU có thể lên tới

  • 16×768=12,288 threads (cho 1st gen., với 16 SM)
  • 30×768=23,040 thread (cho 2nd gen. với 30 SM)

Các threads thuộc cùng 1 warp thì bắt đầu cùng lúc, cùng thực hiện một lệnh, tuy nhiên chúng có thể hoạt động độc lập, i.e. rẽ nhánh (branching) có thể xảy ra. Nếu branching xảy ra, ví dụ theo 2 hướng A và B, thì một số threads sẽ thực hiện A trước, và các threads còn lại thực hiện B sẽ đợi cho tới khi A xong sẽ thực hiện. Chính vì thế, chúng có thể được đồng bộ với nhau, ví dụ: đảm bảo là mọi threads đều đã đọc xong dữ liệu của nó trước khi cho phép mỗi thread thực hiện tiếp tác vụ, hay là các threads cần phải kết thúc trước khi trở lại chương trình gọi.

Dù các GPGPU đời mới vẫn cho phép tính toán với double-precision point, GPGPU mặc định thực hiện tính toán với single-precision point. Do đó, nếu muốn dùng độ chính xác kép, ta cần chỉ định lúc biên dịch với compiler flag–gpu-name sm_13” cho ngôn ngữ C (với Fortran sẽ đề cập sau). Cần nhắc lại là mỗi SM (ví dụ: gồm 8 SP) chỉ có 1 ALU cho tính toán độ chính xác kép. Vì thế, tốc độ lúc dùng double-precision point sẽ bị chậm lại (ví dụ: 8 lần).

SIMT_Cuda

1 dòng lệnh, nhưng chạy đa luồng

Khi lập trình, cần chỉ rõ grid và block dimensions (kích thước tối đa là tùy vào giới hạn phần cứng). Mỗi blockthread đều có một số định danh (ID). Để xác định duy nhất 1 thread, ta cần biết thread block ID, và 1 local thread ID.

Tối ưu phân chia công việc

Để chương trình chạy tốt, ta cần cân bằng giữa số blocks và số thread trong mỗi block khi gọi một kernel. Nếu số threads lớn, thì thì số registers available cho mỗi thread sẽ giảm đi, nghĩa là khả năng lưu dữ liệu truy xuất nhanh bị giảm lại. Mỗi SM có 8,192 registers; mỗi thanh ghi 16KB. Nếu dùng tối đa 768 threads, thì mỗi thread chỉ có thể có khoảng 10 registers, hay ~ 160KB cho bộ nhớ nhanh.

Chú ý là mỗi SM có tối đa 768 threads song song, nên nếu blocks gồm 512 threads thì chỉ có thể có 1 block active trên 1 SM. Như vậy sẽ phí đi 768-512=256 threads. NVIDIA khuyên nên dùng blocks có kích thước trong khoảng từ 128 (1×128, 8×16, 8x8x2, 4x4x8) đến 256 thread (1×256, 16×16, 8x8x16).

Đồng bộ hoạt động giữa các threads:

Các threads của một kernel chạy bất đồng bộ với nhau.  Để đảm bảo sự đồng bộ giữa chúng, ví dụ, khi cùng truy cập một dữ liệu dùng chung, ta dùng cudaThreadSynchronize().

Lúc đó, chương trình sẽ chỉ thực thi kernel kế tiếp khi mà mọi threads của kernel trước đã kết thúc.

Tổng kết

CUDA cung cấp:

  • mô hình về bộ nhớ
  • mô hình về lập trình

CUDA cho phép:

  • codes có thể chạy trên mọi qui mô (scalable or “extensible” parallel programming model), i.e. không cần biên dịch lại khi chạy trên máy có số processor khác. Cụ thể như sau, nếu ta cần chạy 1 kernel với grid gồm 8 blocks. Nếu GPU có 2 core, thì nó sẽ là song song từng 2 blocks, nếu nó có 4 cores, thì sẽ tự động song song từng 4 blocks

CUDA_scalable

  • hỗ trợ codes chạy trên môi trường host lẫn device.
  • mỗi thời điểm chỉ 1 kernel được thực thi, tuy nhiên nhiều threads (arrays of threads) đồng thời thực thi kernel đó và trên vài trăm cores (thread processors), i.e. tương ứng vài ngàn threads song song.
  • các threads trong cùng 1 block có thể phối hợp với nhau (ví dụ: chia sẻ dữ liệu, đồng bồ về thực thi)
  • nhà lập trình chỉ định số lượng threads/blocks, còn phần cứng sẽ quản lí thread processors nào xử lí chúng
  • chính cơ chế thread blocks (với 1-, 2-, và 3-D), và grid of blocks (với 1-, 2-D) cho phép đơn giản hóa việc xử lí các dữ liệu 2 chiều và 3 chiếu (ví dụ: xử lí ảnh, tính toán phương trình vi phần từng phần (PDE ) …)
Mỗi kernel có thể được thực thi song song

Mỗi kernel có thể được thực thi song song

Sự khác biệt giữa GPU và CPU threads:

  • GPU thread là cực kì nhỏ gọn (lightweight)
  • Mỗi GPU có thể chạy hàng ngàn threads song song (multi-core CPU chỉ có thể chạy vài chục threads song song)
  • Threads tự động kết thúc sau khi xong

Kiến trúc tương lai

Có thể nói thêm đôi chút về viễn cảnh của tương lai khi mà CPU+GPU có thể được kết hợp lại (CPU/GPU fusion). Một hướng khác là sự tích hợp giữa CPU core và accelerator core, gọi là APU – Accelerated Processing Unit, theo AMD. Nó có thể là một sự kết hợp giữa CPU cores (scalar processing cores), GPU cores (parallel processing cores), và fixed-function accelerator cores. Intel thì đang phát triển kiến trúc mới gọi là Larrabee còn IBM, Sony và Toshiba liên kết để tạo ra kiến trúc  CELL.

Tham khảo:

  1. https://computing.llnl.gov/tutorials/parallel_comp/
  2. http://www.anandtech.com/video/showdoc.aspx?i=3651
  3. http://en.wikipedia.org/wiki/Comparison_of_MPI,_OpenMP,_and_Stream_Processing
  4. http://www.cs.utexas.edu/~skeckler/wild04/Paper14.pdf (Stream Processing in General-Purpose Processors)
  5. http://saahpc.ncsa.illinois.edu/papers/Chai_paper.pdf
  6. https://visualization.hpc.mil/wiki/GPGPU
  7. http://www.cs.tut.fi/~pk/seminaarit/gpu-2009/cuda-esitys.pdf
  8. http://www.sdsc.edu/us/training/assets/docs/NVIDIA-02-BasicsOfCUDA.pdf
  9. http://www.computerpoweruser.com/editorial/article.asp?article=articles/archive/c0807/30c07/30c07.asp&guid=
  10. http://www.tomshardware.com/reviews/nvidia-cuda-gpu,1954-7.html (warp)
  11. Larrabee: A Many-Core x86 Architecture for Visual Computing
  12. http://en.wikipedia.org/wiki/Larrabee_(GPU)
  13. http://en.wikipedia.org/wiki/CELL

Written by vietnamen

Tháng Mười 3, 2009 at 11:06 chiều

Phần 0: From GPU to GPGPU

leave a comment »

Các máy tính đời đầu không có màn hình. Sau này, khi màn hình ra đời, nhu cầu xử lí hình ảnh, video, đòi hỏi phải có một bộ xứ lí riêng, gọi là GPU, được gắn trên các card màn hình.
Graphic Processing Unit (GPU), ngày nay, đã có thể xử lí mạnh hơn các bộ vi xử lí trung tâm (Central Processing Unit – CPU) trong việc xử lí tính toán. Vai trò của GPU đã vượt ra ngoài phạm vi truyền thống với khái niệm mới General-Purpose computation on GPU (GPGPU) có thể phục vụ cho các ứng dụng cần tính toán nhiều, chứ không đơn thuần cho đồ họa.

GPGPU đầu tiên của AMD công bố (Nov. 2006) với tên gọi CTM (Close to Metal), về sau đổi tên là Stream. GPGPU đầu tiên của NVIDIA là nền tảng Geforces 8 series. Để cung cấp khả năng truy xuất các tập lệnh của các GPU này, NVIDIA cung cấp một Software Development Kit (SDK) và các Application Programming Interface (API) đi kèm viết cho ngôn ngữ lập trình C gọi là Compute Unified Device Architecture (CUDA). Các GPGPU của NVIDIA còn gọi là CUDA-capable GPU. Hãng AMD cũng đưa ra một SDK/API tương tự  cho các GPGPU dòng ATI của hãng, gọi là Brook+.

Gần đây, PGI cũng cho ra phiên bản Fortran compiler với CUDA-support. Với CUDA, các programmers có thể dùng ngôn ngữ C để viết các chương trình và dùng nvcc cùng gcc để biên dịch nó. Chương trình được biên dịch có thể sử dụng khả năng tính toán của GPGPU.

Tuy nhiên, vì các tập lệnh của các GPU được chuyên biệt hóa phục vụ cho đồ họa, phạm vi ứng dụng cho các ứng dụng tính toán cũng bị giới hạn. Bù lại, khả năng xử lí song song của GPU là cực kì mạnh. Trước đây, số lệnh được thực thi tròng vòng 1s (instructions per seconds) là đơn vị dùng để đo tốc độ của CPU. Ngày nay, trong lĩnh vực tính toán khoa học, với việc tập trung nhiều vào khả năng tính toán (có sử dụng dấu chấm động – floating-point calculations), đơn vị để đo lường sức mạnh tính toán là floating point operations per second (FLOPS). LINPACK benchmarks là công cụ thường dùng để đo lường hiệu năng theo FLOPS.

Sau đây là một vài so sánh

  • Tháng 5, 2008, một siêu máy tính của quân đội Mĩ do IBM chế tạo đạt tốc độ 1petaFLOPS = 1000 petaFLOPS ~ 1.026 quadrillion calculations per second
  • Dự kiến, có thể tạo ra máy tính đạt tốc độ 10 petaFLOPS năm 2012
  • Tháng 6, 2008, GPGPU của AMD (ATI Radeon HD4800 series) đạt tốc độ 1teraFLOPS = 1000 gigaFLOPS
  • Năm 2008, bộ vi xử lí quad-core của Intel (Intel Core i7 965 XE) chỉ mới đạt 70 gigaFLOPS = 10^9 FLOPS – độ chính xác kép
  • Bộ vi xử lí SPARC64 VIIIfx Venus CPU của Fujitsu (năm 2008)  đạt tốc độ 128 GigaFLOPS.

Hạn chế hiện nay của GPGPU (chip G80 cho NVIDIA, chip … cho AMD) là nó chỉ hiệu quả với các phép toán độ chính xác đơn.

  1. với single precision calculations
  • nVidia’s Tesla C1060 GPU: 933 GFLOPS
  • AMD’s FireStream 9270: 1200 GFLOPS
  1. với double precision calculations
  • Nvidia Tesla C1060: 78 GFLOPS
  • AMD FireStream 9270: 240 GFLOPS

Hiện nay, CUDA được ưa chuộng hơn vì ngôn ngữ dùng là C/C++ thay vì dùng ngôn ngữ cấp thấp (assembly) như đối với CMI. Vì thế, các bài viết của tôi sẽ tập trung vào công nghệ CUDA.

Tăng trưởng về tốc độ so sánh giữa CPU và GPU

Tăng trưởng về tốc độ so sánh giữa CPU và GPU

Các thế hệ tiếp theo của CUDA-capable GPU là NVIDIA 9 series, 100 series, 200 series, Quadro. Tuy nhiên, phải đợi đến khi NVIDIA Tesla ra đời mới có 1 GPGPU đúng nghĩa. Đây là dòng GPU không có ngõ ra video, do đó, nó được thiết kế thuần túy cho tính toán. Các models hiện nay của Tesla là C870, D870, S870, C1060, S1070; đi theo 3 dạng: dạng card, dạng “external box” và dạng rack-mount server form, i.e. “C” viết tắt của “Card”, “D” viết tắt của “Desktop” và “S” viết tắt của “Server”. Các Tesla card thế hệ đầu (T8) sau dùng chip G80 (các dòng giá rẻ như G84, G86 không được sử dụng)

  1. C870 dùng chuẩn PCI Express x16 với thông số 1.5GB bộ nhớ và sức mạnh 518 GigaFLOPS – dùng chip G80 với 128 cores. Điện năng: 170W. (~1,500$)
  2. D870 được thiết kế dưới dạng 1 thùng độc lập (external box) gồm 2 card C870 nối với máy tính qua 1 cable PCI Express mở rộng, với sức mạnh 1 TetraFLOPS. Điện năng: 550W. (~7,500$)
  3. S870 được thiết kết với 4 card C870 theo kiến trúc Rack-Mount server form (1U), có thể nối với máy chủ qua cable PCI Express mở rộng, với sức mạnh đến 2 TetraFLOPS. Điện năng: 800W. (~12,000$)

CHÚ Ý: Các thông số trên là cho single-precision calculations.

Khả năng tiêu hao điện thấp là một điểm mạnh của GPGPU. Một GeForce GPU có khả năng tương đương 40 vi xử lí x86. Trong khi, với số lượng vi xử lí x86 đó cần 1600W.

Tesla D870 (price: 7500 USD)

Tesla S870 (price: 12000 USD)
Tesla S870 (price: 12000 USD)

Tesla T10 (thế hệ 2) với 2 phiên bản C1060 và S1070. T10 sử dụng chip GT200 series gồm 240 cores (gần gấp đôi so với 128 cores của G80 series) có khả năng xử lí 1 teraFLOPS gần gấp đôi Tesla GPU thế hệ đầu, cùng với độ chính xác cao hơn (64bit) và RAM tăng. Một số hình ảnh về GPU Tesla có thể tham khảo ở đây.

  1. C1060 giống như 1 card màn hình hơi to, kết nối cổng PCI Express x16, chứa 1 chip T10P (mỗi core tốc độ 1.33GHz, 4GB 800MHz GDDR3 RAM, băng thông bộ nhớ là 102GB/s), tốc độ 900 gigaFLOPS. Điện năng: 160W. ($1699)
  2. S1070 1U kích thước nhỏ hơn, vơi 4 chip T10P, i.e. 4×240=960 cores, tốc độ 4 tetraFLOPS. Mỗi chip tốc độ 1.5GHz, 16GB 800MHz GDDR3 RAM với băng thông bộ nhớ là 408 GB/s. Điện năng: 700W. ($7995 – 2PCIe interfaces, $8295 – 1 PCIe interface)

Tesla Fermi (thế hệ 3) dự kiến sẽ ra đời sớm. Nó dùng chip G300 series, GT300 hoặc NV60, không có ngõ ra DVI-I. Chip thế hệ mới có 512 cores (hơn gấp đôi so với thế hệ 2) hỗ trợ chuẩn IEEE 754-2008 floating-point standard, công nghệ sữa lỗi ECC (error-correcting code). Các kernels (khái niệm mới này sẽ được giới thiệu ở các bài sau) có thể chạy song song. Bên cạnh đó là sự ra đời của Nexus – môi trường để phát triển các ứng dụng cho tính toán, được tích hợp với Microsoft Visual Studio. Sau đây là một số so sánh với 2 thế hệ đầu.

So sánh các thế hệ của Tesla

So sánh các thế hệ của Tesla

Yêu cầu:

  1. máy tính có cài đặt 1 GPGPU: Geforces 8 series, 9 series, 100 series, 200 series, Quadro, and Tesla.
  2. CUDA toolkit, CUDA SDK, NVIDIA driver…
  3. chương trình được viết và biên dịch với compiler (gcc, hay PGI Fortran)

Update:

  1. Intel Nehalem với 4 cores, tốc độ 50GFLOPS, tiêu thụ 80W/h (price: $500).
  2. NVIDIA Fermi chips ~8 times faster than Tesla in double-precision computing, providing ECC (error correcting code) memory, double-precision

 

Tham khảo:

  1. http://gpgpu.org/
  2. http://ati.amd.com/technology/streamcomputing/gpgpu_history.html (GPGPU history from AMD)
  3. http://gpgpu.org/oldsite/data/history.shtml
  4. http://www.techradar.com/news/computing-components/processors/world-s-fastest-cpu-clocked-at-128-gigaflops-599433
  5. http://www.thegioimaychu.vn/forum/showthread.php?p=1127
  6. http://www.thongtincongnghe.com/article/356
  7. http://www.tomshardware.com/reviews/nvidia-cuda-gpu,1954-2.html
  8. http://www.vizworld.com/2009/06/nvidia-and-pgi-partner-to-bring-cuda-to-fortran/
  9. http://www.tgdaily.com/index.php?option=com_content&task=blogsection&id=18&Itemid=41&slideshow=20070620
  10. http://penstarsys.com/previews/graphics/nvidia/tesla/tesla_2.htm
  11. http://www.tomshardware.com/news/Tesla-C1060-S1070,5672.html
  12. http://www.xbitlabs.com/news/video/display/20090930223729_First_Images_of_Nvidia_GeForce_Fermi_Show_Up.html
  13. http://hothardware.com/News/NVIDIA-Unveils-Next-Generation-Fermi-GPU-Architecture/

Written by vietnamen

Tháng Mười 1, 2009 at 3:27 chiều