Vietnamen’s Weblog

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

Archive for Tháng Mười 2009

Công cụ dạy học và nghiên cứu

leave a comment »

Secondary school

Quantum Science Across Disciplines (QSAD):

http://qsad.bu.edu

Phần mềm giúp học sinh hiểu rõ hơn về khoa học lượng tử, với cách nhìn của một học sinh cấp 2-3.

Physics

ABINIT

http://en.wikipedia.org/wiki/ABINIT

Phần mềm

  1. tính tổng năng lượng, mật độ điện tích, cấu trúc electron của 1 hệ thống tạo bởi các electron và hạt nhân dùng DFT (density function theory), pseudopotential…
  2. tối ưu hoá cấu trúc hình học (3D) dựa vào DFT forces, stresses
  3. MD simulation

Chemoinformatics & Bioinformatics

Chemistry Development Kit (SDK)

http://sourceforge.net/apps/mediawiki/cdk/index.php?title=Main_Page

A Java library

  1. 2D structure diagram editor, layout
  2. 3D rendering
  3. support various format (Chemical Markup Language, SMILES, MDL, InChl)
  4. modeling
  5. PDB file (active site detection, sequence to connectivity table)
  6. BioJava interface

Discrete Fourier transform (DFT)

FFTW

http://www.fftw.org

C subroutine library

Written by vietnamen

Tháng Mười 21, 2009 at 1:35 chiều

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

Sự lựa chọn một màn hình LCD tốt

leave a comment »

Để có một màn hình LCD vừa ý, nó không chỉ đơn thuần là lựa chọn thương hiệu mà chính là những thông số kĩ thuật giúp ta đánh giá chính xác hơn chất lượng của một màn hình LCD (Liquid Crystal Display). Có 2 loại màn hình LCS: Passive Matrix và Active Matrix (loại sau phổ biến hơn vì tốc độ xử lí cao). Và Active Matrix dùng TFT (thin film resistor). Sau đây là một vài thông tin thu thập được

  1. Kiểu đầu vào: bắt buộc phải hỗ trợ đầu cắm DVI, chứ không phải kiểu đầu nối 15-chân D-sub. Điều này giúp bạn kết nối với các thế hệ card màn hình mới với chất lượng đồ họa cao hơn.
  2. Độ phân giải, tỉ lệ kích thước: với tỉ lệ kích thước (chiều rộng với chiều cao) là 4:3 thì ta có các độ phân giải 1600×1200, 1280×960, nếu tỉ lệ kích thước là 5:4 thì ta có độ phân giải 1280×1024.
  3. Thời gian đáp ứng: thông số này không quan trọng lắm vì bạn không thấy rõ sự khác biệt.
  4. Độ sâu màu: 6-bit (), 8-bit (16.2M màu) hay 24-bit.
  5. Độ sáng: đơn vị là ‘cd/m2’ (candela per meter squared) hay ‘nits’.
  6. Góc nhìn: với màn hình LCD mới thì thông số này không quan trọng lắm
  7. Tỉ lệ tương phản: nếu độ sáng của phần màn hình hiển thị đen là 0.5 cd/m2, và độ sáng của phần màn hình hiển thị trắng là 250 cd/m2, thì tỉ lệ tương phản 500:1. Tuy nhiên, do cách tính toán của các hãng cũng khác nhau, nên số liệu này cũng không chuẩn xác lắm.

Tổng kết lại, không nên đầu tư vào những chức năng mà mình ít dùng tới.

Thông số mẫu:

Sony SDM-S94
LCD 19″ SXGA LCD (Active Matrix)
pixel pitch: 0.294mm
Anti-glare coating
Scanning Frequency Horizontal: 28-80kHz
Vertical: 45-75Hz
Response Time 25ms (Typical)
Contrast Ratio 600:1 (Typical)
Compatibility 1280 x 1024 (Native)
Brightness 250 cd/m2
Viewing Angle 170 / 170 (Horizontal / Vertical)
Power Working: 50W
Standby/Off: 1.2W
Warranty 3 years parts and labor
Interface DVI
15-pin D-sub

Tham khảo:

  1. http://www.anandtech.com/displays/showdoc.aspx?i=2289&p=2
  2. http://www.anandtech.com/displays/showdoc.aspx?i=2289&p=10

Written by vietnamen

Tháng Mười 5, 2009 at 7:08 chiều

Posted in Linh tinh, Tin học

Tagged with , ,

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