Vietnamen’s Weblog

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

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 lúc 10:07 chiều

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: