Vietnamen’s Weblog

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

Posts Tagged ‘Fortran

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

Bước đầu với Fortran

leave a comment »

Written by vietnamen

Tháng Mười Hai 11, 2008 at 3:44 chiều

Posted in Linh tinh

Tagged with ,