Vietnamen’s Weblog

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

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 lúc 11:37 sáng

Gửi phản hồi

Mời bạn điền thông tin vào ô dưới đây hoặc kích vào một biểu tượng để đăng nhập:

WordPress.com Logo

Bạn đang bình luận bằng tài khoản WordPress.com Log Out / Thay đổi )

Twitter picture

Bạn đang bình luận bằng tài khoản Twitter Log Out / Thay đổi )

Facebook photo

Bạn đang bình luận bằng tài khoản Facebook Log Out / Thay đổi )

Google+ photo

Bạn đang bình luận bằng tài khoản Google+ Log Out / Thay đổi )

Connecting to %s

%d bloggers like this: