Chương 7: Học thông qua thực hành! Lập trình OpenCL

Estimated reading: 38 minutes 17 views

Truyền nhận dữ liệu giữa Host (Máy chủ) và Kernel (Hạt nhân) /

Xử lý đa luồng /

Các điều kiện ràng buộc của Kernel.

_____________________

Kể từ bài viết này, chúng ta sẽ sử dụng chương trình chap4/opencl-hello-error-check đã được tạo ra ở Chương 4 làm nền tảng, sau đó sẽ tiếp tục học lập trình GPU bằng cách thêm thắt các tính năng mới vào đó. Tại những chỗ có thay đổi so với chương trình gốc, tôi đã đánh dấu bằng ký hiệu //☆.

Trong bài viết này, tôi sẽ giải thích 3 khái niệm cơ bản sau:

・Truyền nhận dữ liệu giữa Máy chủ (Host / ホスト) và Hạt nhân (Kernel / カーネル)

・Xử lý Đa luồng (Multithreading / マルチスレッド)

・Các Điều kiện ràng buộc (Constraints / 制約事項) của Hạt nhân (Kernel / カーネル) OpenCL

1. Truyền nhận dữ liệu giữa Máy chủ (Host) và Hạt nhân (Kernel)

Cách truyền tham số (Argument / 引数) vào hàm hạt nhân

chap5-1/opencl-arg-test

● Hàm hạt nhân (Kernel function / カーネル関数)

Chúng ta sẽ thêm quá trình xử lý truyền tham số (Argument / 引数) vào hàm hạt nhân. Hàm hạt nhân được trình bày ở Danh sách 1(a) nhận 2 tham số là biến kiểu int tên a và b, sau đó thực hiện tính toán $c = a + b$ và hiển thị kết quả.

Danh sách 1: Đoạn mã truyền tham số (Argument) từ Mã máy chủ (Host code) vào Hàm hạt nhân (chap5-1/opencl-arg-test)

				
					// (a) Hàm hạt nhân (Kernel function)
__kernel void arg_test(int a, int b)
{
  int c;
  
  c = a + b;
  printf("[Kernel] %d + %d = %d\n", a, b, c);
}


// (b) Mã bổ sung phía Mã máy chủ (Host code)
cl_int ha = 10, hb = 20;

// Thiết lập tham số cho Kernel
clSetKernelArg(kernel, 0, sizeof(cl_int), (void *)&ha);
clSetKernelArg(kernel, 1, sizeof(cl_int), (void *)&hb);

// Thực thi Kernel
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, gws, lws, 0, NULL, NULL);

● Phía Mã máy chủ (Host code)… Thiết lập giá trị tham số trước khi thực thi Kernel

Ở phía Mã máy chủ (Host code / ホスト・コード) chịu trách nhiệm gọi đoạn mã hạt nhân này, tôi đã chèn thêm đoạn xử lý của Danh sách 1(b). Trước khi thực thi Kernel, bằng cách gọi hàm clSetKernelArg(), chúng ta sẽ thiết lập giá trị của các tham số (Argument / 引数) muốn truyền vào cho hàm hạt nhân. Mặc dù có cảm giác hơi phiền phức khi cứ mỗi lần tăng thêm 1 tham số thì lại phải tốn thêm 1 lần gọi API, nhưng nếu bạn muốn truyền tham số theo dạng truyền giá trị (Pass-by-value) thì đoạn mã sẽ có hình hài như vậy. Ngoài phương pháp này, còn có cách truyền lượng lớn dữ liệu thông qua Con trỏ (Pointer / ポインタ) sẽ được tôi giải thích ở phần sau. Tham số thứ 2 của hàm clSetKernelArg() là thông số (Parameter) chỉ định xem đó là tham số thứ mấy. Ngoài ra, tham số thứ 4 – nơi chứa giá trị muốn truyền cho Kernel – bắt buộc phải là một Con trỏ (Pointer / ポインタ). Do đó, nếu muốn truyền giá trị gì đó cho Kernel, bạn phải lưu trữ giá trị đó vào một biến từ trước giống như trong đoạn mã mẫu.

● Kết quả thực thi

Kết quả chạy chương trình được hiển thị bên dưới.

Bash

$ ./main
  :
[Kernel] 10 + 20 = 30

Giá trị của các biến được thiết lập bằng lệnh clSetKernelArg() đã được truyền vào các tham số (Argument / 引数) của hàm hạt nhân một cách chính xác.

Cách trả về kết quả xử lý từ Hàm hạt nhân

chap5-1/opencl-retval-test

● Hàm hạt nhân… Cấp phát Khu vực dữ liệu (Data area / データ領域) trên VRAM

Các hàm hạt nhân không thể trả về giá trị (bằng lệnh return). Nếu cần trả về bất kỳ kết quả xử lý nào từ Hàm hạt nhân, bạn bắt buộc phải dùng tham số (Argument / 引数).

Như ở Danh sách 2(a), tôi đã thêm một Con trỏ (Pointer / ポインタ) trỏ tới biến c làm tham số dùng để trả về giá trị cho hàm hạt nhân của Danh sách 1(a).

Phương pháp truyền một con trỏ của biến khi muốn trả về giá trị thông qua tham số là một kỹ thuật quá sức quen thuộc trong ngôn ngữ C, nên chắc hẳn sẽ không có gì lạ lẫm. Tuy nhiên, do bộ nhớ ở phía Máy chủ (Host / ホスト) (tức CPU) và bộ nhớ ở phía GPU (VRAM) được xử lý riêng biệt, hàm hạt nhân không tài nào truy cập trực tiếp vào bộ nhớ của Máy chủ được. Do đó, Khu vực dữ liệu (Data area / データ領域) dùng để lưu trữ biến c đó bắt buộc phải được cấp phát ở trên VRAM. Chính vì vậy, tôi đã thêm từ khóa bổ nghĩa (Modifier) __global cho tham số c nhằm báo cho Trình biên dịch C (C Compiler / Cコンパイラ) biết rằng đích đến mà con trỏ biến c đang trỏ tới là nằm trên VRAM.

● Phía Mã máy chủ (Host code)… Copy kết quả tính toán trên VRAM về Máy chủ

Các đoạn thay đổi ở phía Mã máy chủ được thể hiện trong Danh sách 2(b).

Sử dụng Hình 1 làm tham chiếu, tôi sẽ giải thích hoạt động của chương trình opencl-retval-test. Trong Mã máy chủ, bằng cách gọi clCreateBuffer(), một khu vực dữ liệu (Data area / データ領域) dành riêng cho biến c sẽ được cấp phát trên VRAM (Bước ① trong Hình 1). Giá trị trả về mem_obj của hàm clCreateBuffer() hoạt động giống như một Handle (Định danh) dùng để phân biệt vùng bộ nhớ đã được cấp phát. Tiếp theo, ta gọi clSetKernelArg() để truyền con trỏ của vùng bộ nhớ đã cấp phát trên VRAM vào làm tham số thứ 3 cho hàm hạt nhân (Bước ④ trong Hình 1).

Hình 1: Luồng trả về kết quả xử lý từ GPU về cho Máy chủ (Host)

Quá trình của chương trình opencl-retval-test

Danh sách 2: Đoạn mã trả về kết quả xử lý từ hàm hạt nhân (chap5-1/opencl-retval-test)

				
					// (a) Hàm hạt nhân (Kernel function)
__kernel void retval_test(int a, int b, __global int *c)
{
  *c = a + b;
  printf("[Kernel] %d + %d = %d\n", a, b, *c);
}

// (b) Đoạn mã thay đổi phía Mã máy chủ (Host code)
cl_int ha = 10, hb = 20, hc;
cl_mem mem_obj;

// Cấp phát thiết bị bộ nhớ (VRAM) cần thiết
mem_obj = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(hc), NULL, &ret);

// Thiết lập tham số cho Kernel
clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&mem_obj);

// Thực thi Kernel

// Copy dữ liệu đầu ra từ VRAM về bộ nhớ Host (Máy chủ)
clEnqueueReadBuffer(queue, mem_obj, CL_TRUE, 0, sizeof(hc), &hc, 0, NULL, NULL);

// In kết quả tính toán trên GPU
printf("[Host] hc=%d\n", hc);

// Giải phóng các tài nguyên (Resource)
ret = clReleaseMemObject(mem_obj);

Khi Kernel thực thi (Bước ⑤ trong Hình 1), vì phía hàm hạt nhân [Danh sách 2(a)] đang chạy dòng lệnh *c = a + b;, giá trị của $a+b$ sẽ được ghi vào khu vực dữ liệu dành cho biến c trên VRAM. Mặc dù kết quả tính toán đã hiện diện trên VRAM, nhưng CPU Máy chủ không thể với tay tới nó. Thế nên, bằng cách gọi clEnqueueReadBuffer(), dữ liệu đầu ra của GPU đang chễm chệ trên VRAM (mem_obj) sẽ được sao chép về biến hc ở phía Máy chủ (Bước ⑥ trong Hình 1).

● Kết quả thực thi

Kết quả chạy của chương trình opencl-retval-test được hiển thị bên dưới.

Bash

$ ./main
  :
[Kernel] 10 + 20 = 30
[Host] hc=30

Kết quả của quá trình xử lý tính toán trên GPU đã được truyền thành công về phía CPU Máy chủ (Host / ホスト).

Truyền nhận lượng lớn dữ liệu giữa Máy chủ (Host) và Hạt nhân (Kernel)

chap5-1/opencl-add-uint

● Hàm hạt nhân… Đổi cách truyền tham số a, b sang dùng Con trỏ (Pointer)

Đến tận bây giờ, những gì chúng ta truyền qua lại giữa Mã máy chủ và Mã hạt nhân (Kernel code / カーネル・コード) chỉ là các biến chứa Giá trị vô hướng (Scalar value / スカラ値). Tuy nhiên, thông thường mục đích ta nhờ vả GPU là để tăng tốc (Acceleration / 高速化) xử lý tính toán trên một tập dữ liệu lớn. Do đó, việc hiện thực hóa truyền nhận lượng lớn dữ liệu giữa Máy chủ và Hạt nhân là điều bắt buộc.

Trong hàm hạt nhân ở Danh sách 3(a), phép tính $c=a+b$ đang được thực hiện trên 8 khối dữ liệu. Để truyền một lượng dữ liệu lớn (mặc dù trong ví dụ này chỉ là 2 mảng (Array / 配列) có 4 phần tử), tôi đã sửa đổi đoạn mã so với Danh sách 2(a) sao cho các tham số a, b cũng được truyền vào dưới dạng Con trỏ (Pointer / ポインタ) trỏ tới VRAM.

● Phía Mã máy chủ… Cấp phát vùng nhớ trên VRAM cho cả biến a, b

Những thay đổi ở phía Mã máy chủ được thể hiện trong Danh sách 3(b).

Trong Mã máy chủ, tôi đã thêm các lệnh gọi clCreateBuffer() để cấp phát thêm khu vực dữ liệu (Data area / データ領域) dành cho biến a và b trên VRAM.

Tôi xin giải thích một chút về các thông số (Parameter) được truyền vào làm tham số thứ 2 của hàm clCreateBuffer(). Đối với biến a, b, do GPU chỉ thực hiện thao tác Đọc (Read), nên khi cấp phát mem_obj_a và mem_obj_b, chúng ta sẽ nhét cờ CL_MEM_READ_ONLY vào. Tương tự đối với biến c, do GPU chỉ Ghi (Write) dữ liệu, ta sẽ truyền cờ CL_MEM_WRITE_ONLY.

Trước khi bắt đầu cho Kernel thực thi, quá trình copy dữ liệu từ các biến ha, hb phía Máy chủ sang mem_obj_a, mem_obj_b trên VRAM sẽ được tiến hành bằng cách gọi hàm clEnqueueWriteBuffer().

Danh sách 3: Đoạn mã truyền lượng lớn dữ liệu giữa Máy chủ (Host) và Hạt nhân (Kernel) (chap5-1/opencl-add-uint)

				
					// (a) Hàm hạt nhân (Kernel function)
__kernel void much_data(__global uint *a, __global uint *b, __global uint *c)
{
  *c++ = *a++ + *b++;
  *c++ = *a++ + *b++;
  *c++ = *a++ + *b++;
  *c++ = *a++ + *b++;
}

// (b) Đoạn mã thay đổi phía Mã máy chủ (Host code)
// Số lượng dữ liệu mục tiêu cần tính toán
#define N_DATA    4

int main(int argc, char *argv[])
{
  cl_uint ha[N_DATA] = { 1, 2, 3, 4 };
  cl_uint hb[N_DATA] = { 20, 30, 40, 50 };
  cl_uint hc[N_DATA];

  // Cấp phát thiết bị bộ nhớ (VRAM) cần thiết
  mem_obj_a = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(ha), NULL, &ret);
  mem_obj_b = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(hb), NULL, &ret);
  mem_obj_c = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(hc), NULL, &ret);

  // Copy dữ liệu đầu vào từ bộ nhớ Host (Máy chủ) sang VRAM
  clEnqueueWriteBuffer(queue, mem_obj_a, CL_TRUE, 0, sizeof(ha), ha, 0, NULL, NULL);
  clEnqueueWriteBuffer(queue, mem_obj_b, CL_TRUE, 0, sizeof(hb), hb, 0, NULL, NULL);

  // Thiết lập tham số cho Kernel
  clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&mem_obj_a);
  clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&mem_obj_b);
  clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&mem_obj_c);

  // Thực thi Kernel

  // Copy dữ liệu đầu ra từ VRAM về bộ nhớ Host (Máy chủ)
  clEnqueueReadBuffer(queue, mem_obj_c, CL_TRUE, 0, sizeof(hc), hc, 0, NULL, NULL);

  // In kết quả tính toán trên GPU
  for (i = 0; i < N_DATA; i++) {
    printf("[Host] hc[%d]=%d\n", i, hc[i]);
  }
}

● Kết quả thực thi

Kết quả chạy của chương trình opencl-add-uint được hiển thị bên dưới.

Bash

$ ./main
  :
[Host] hc[0]=21
[Host] hc[1]=32
[Host] hc[2]=43
[Host] hc[3]=54

Các xử lý tính toán trên tập dữ liệu lớn đã được phía GPU nhai ngấu nghiến một cách chính xác.

＊ ＊ ＊

Các phương thức trao đổi thông tin giữa Máy chủ và Hạt nhân được tổng kết lại ở Hình 2. Thứ duy nhất có thể tuồn từ Mã máy chủ sang Mã hạt nhân (Kernel code / カーネル・コード) qua đường tham số (Argument / 引数) chỉ là những Giá trị vô hướng (Scalar value / スカラ値) như kiểu biến int a, hoặc các Con trỏ (Pointer / ポインタ) trỏ tới VRAM. Thêm vào đó, để có thể xử lý tính toán trên GPU, bạn buộc phải sử dụng Bộ đệm bộ nhớ (Memory buffer / メモリ・バッファ) đã được cấp phát trên VRAM, và chỉ duy nhất Mã hạt nhân mới có quyền thọc tay vào Bộ đệm bộ nhớ trên VRAM này.

Hình 2: Phương thức truyền đạt thông tin giữa Máy chủ – Hạt nhân (Host – Kernel / ホスト – カーネル)

Những thứ có thể truyền sang Hạt nhân chỉ có thể là Giá trị vô hướng hoặc Con trỏ trên VRAM.

                         [ Mã máy chủ ]
                               |
       |-----------------------------------------------|
       |                                               |
 [ clSetKernelArg() ]                     ( Có thể tự do truy cập bộ nhớ )
       |                                               |
       v                                               v
   [ Tham số ]                                 [ Bộ đệm bộ nhớ ]

---------------------------------------------------------------------- Bộ nhớ Máy chủ
---------------------------------------------------------------------- VRAM

       |                               |               ^
       |                    clEnqueueWriteBuffer()  clEnqueueReadBuffer()
       |                               |               |
       v                               v               |
   [ Tham số ]                                 [ Bộ đệm bộ nhớ ]
       |                                               |
       |                                  ( Có thể tự do truy cập bộ nhớ )
       v                                               v
                         [ Mã hạt nhân ]


    《 Những thứ có thể truyền qua tham số 》
    ◯ Giá trị vô hướng
    ◯ Con trỏ trên VRAM
    × Con trỏ trên bộ nhớ Máy chủ

Khi trao đổi dữ liệu đầu vào/đầu ra giữa phía Máy chủ và GPU, bạn bắt buộc phải xài 2 API là clEnqueueWriteBuffer() và clEnqueueReadBuffer() để tiến hành quá trình copy dữ liệu.

2. Xử lý Đa luồng (Multithreading)

Thực hiện xử lý Đa luồng trên GPU (1 Chiều)

chap5-2/opencl-mt-1d

● Tung hết bản lĩnh nhờ xử lý Đa luồng (Multithreading / マルチスレッド)

Trong Danh sách 1 đến Danh sách 3, chúng ta chỉ cho Hàm hạt nhân chạy đúng 1 lần duy nhất. Tuy nhiên, thời khắc GPU thực sự tung hết bản lĩnh là khi một đội quân đông đảo các Phần tử xử lý (Processor element / プロセッサ・エレメント, từ đây viết tắt là PE) tiến hành xử lý Đa luồng (Multithreading / マルチスレッド), tức là thực hiện các phép toán song song đồng thời trên một lượng lớn dữ liệu (Hình 3). Để vắt kiệt hiệu năng tính toán mà GPU sở hữu, việc ứng dụng xử lý Đa luồng này là điều bắt buộc, và nếu bạn xài OpenCL thì việc đa luồng hóa sẽ dễ như trở bàn tay.

[ xxx.c ] ==( Biên dịch )==> [ Mã máy trần CPU ] ==( Thực thi )==> [ CPU Máy chủ ]
(a) Luồng thực thi của Mã máy chủ

                             ( Tất cả các PE đều thực thi cùng một Mã hạt nhân )
                                                          |
                                                          v
[ xxx.cl ] ==( Biên dịch )==> [ Mã máy trần GPU ] ==( Thực thi )==> [ PE ] [ PE ] [ PE ] [ PE ]
                                                                    [ PE ] [ PE ] [ PE ] [ PE ]
(b) Luồng thực thi của Mã hạt nhân

Hình 3: GPGPU phát huy bản lĩnh thực sự chính là nhờ quá trình xử lý Đa luồng (Multithreading / マルチスレッド) thực thi các phép toán song song đồng thời

Trong Mã hạt nhân (Kernel code / カーネル・コード) từ Danh sách 1 đến Danh sách 3, trong số hàng ngàn PE (Phần tử xử lý) bên trong GPU, chỉ có vỏn vẹn 1 cái được đưa vào hoạt động.

● Cách tiến hành xử lý Đa luồng

▶ Phía Mã máy chủ… Chỉ định tổng số luồng (Thread)

Để chạy xử lý Đa luồng (Multithreading / マルチスレッド), ở phía Mã máy chủ, khi gọi hàm clEnqueueNDRangeKernel(), bạn sẽ dùng tham số thứ 5 là gws để ấn định tổng số lượng luồng mà bạn muốn GPU bung ra như Danh sách 4(a). Theo thuật ngữ của OpenCL, số lượng luồng được kích hoạt trên thiết bị mục tiêu (Target device) này được gọi là Kích thước công việc toàn cục (Global work size / グローバル・ワーク・サイズ).

Do tham số thứ 3 của lệnh clEnqueueNDRangeKernel() đang ngầm định rằng cách chỉ định số lượng luồng là theo dạng không gian 1 Chiều (1D), nên trong đoạn mã này chỉ có giá trị được nhồi vào gws[0] là mang ý nghĩa. Nói tóm gọn, trong Danh sách 4(a), Hàm hạt nhân sẽ được thực thi tổng cộng 10 lần ở bên phía GPU.

Danh sách 4: Mã thực hiện xử lý Đa luồng (chap5-2/opencl-mt-1d)

				
					// (a) Phía Mã máy chủ (Host code) chỉ định tổng số luồng
size_t gws[2], lws[2];

// Thực thi Kernel
set_work_size(gws, 10, 0);
set_work_size(lws, 1, 0);
clEnqueueNDRangeKernel(queue, kernel, 1, NULL, gws, lws, 0, NULL, NULL);

//(b) Bên trong Hàm hạt nhân có thể lấy được ID nhận dạng xem bản thân nó là luồng thứ mấy
__kernel void kernel_func(void)
{
  uint x_gid = get_global_id(0);
  
  printf("[Kernel] x_gid=%d\n", x_gid);
}

▶ Phía Hạt nhân (Kernel)… Lấy được thông tin ID của luồng hiện tại

Bên trong Hàm hạt nhân, bản thân cái luồng (Thread) vừa được kích hoạt có thể lấy được thông tin ID để biết xem nó đang là luồng thứ mấy. Trong Hàm hạt nhân ở Danh sách 4(b), bằng cách gọi hàm get_global_id(), hệ thống sẽ lấy được thông tin ID (một số nguyên bắt đầu từ 0) của luồng đã được đánh thức. Giá trị truyền vào làm tham số cho hàm get_global_id() là để chỉ định xem bạn muốn lấy ID của trục (Parameter) thứ mấy trong cấu trúc Kích thước công việc toàn cục (Global work size) đã được cài đặt ở gws. Tức là trong trường hợp này, ứng với các giá trị đã nạp vào mảng gws[], nó sẽ trả về chỉ số (Index) tương ứng với giá trị lưu trong gws[0].

● Kết quả thực thi

Kết quả chạy chương trình opencl-mt-1d được hiển thị bên dưới.

Bash 

$ ./main
  :
[Kernel] x_gid=0
[Kernel] x_gid=1
[Kernel] x_gid=2
[Kernel] x_gid=3
[Kernel] x_gid=4
[Kernel] x_gid=5
[Kernel] x_gid=6
[Kernel] x_gid=7
[Kernel] x_gid=8
[Kernel] x_gid=9

Đúng y như những gì đã cài cắm trong Kích thước công việc toàn cục (Global work size / グローバル・ワーク・サイズ), Hàm hạt nhân đã được gọi réo gọi tổng cộng 10 lần.

Thực hiện xử lý Đa luồng trên GPU (2 Chiều)

chap5-2/opencl-mt-2d

● Cách chỉ định số lượng luồng trong không gian 2 Chiều

▶ Phía Mã máy chủ (Host code)

Tôi sẽ giải thích ví dụ khi muốn chỉ định số lượng luồng trong không gian 2 Chiều. Mã máy chủ chỉ việc thay đổi các giá trị cài đặt trong mảng gws[] và đổi giá trị truyền vào tham số thứ 3 của hàm clEnqueueNDRangeKernel() thành số 2 giống như ở Danh sách 5(a). Trong ví dụ này, tổng lượng luồng đang được định cấu hình bằng 2 thông số trên trục không gian X và Y. Kết quả là bên phía GPU, tổng cộng sẽ có $5 \times 2 = 10$ luồng (Thread) được kích hoạt lên.

Danh sách 5: Mã thực hiện xử lý Đa luồng (chap5-2/opencl-mt-2d)

				
					// (a) Phía Mã máy chủ (Host code)
// Thực thi Kernel
set_work_size(gws, 5, 2);
set_work_size(lws, 1, 1);
clEnqueueNDRangeKernel(queue, kernel, 2, NULL, gws, lws, 0, NULL, NULL);

// (b) Phía Mã hạt nhân (Kernel code)
__kernel void kernel_func(void)
{
  uint x_gid = get_global_id(0);  // ☆
  uint y_gid = get_global_id(1);  // ☆
  
  printf("[Kernel] gid=(%d,%d)\n", x_gid, y_gid); // ☆
}

▶ Hàm hạt nhân (Kernel function)

Mã hạt nhân (Kernel code / カーネル・コード) như được thấy ở Danh sách 5(b) sẽ gọi hàm get_global_id() 2 lần liên tiếp để móc ra cả số ID theo trục x và ID theo trục y rồi đem đi in (print).

● Kết quả thực thi

Kết quả chạy chương trình opencl-mt-2d được hiển thị bên dưới.

Bash

$ ./main
  :
[Kernel] gid=(0,0)
[Kernel] gid=(1,0)
[Kernel] gid=(2,0)
[Kernel] gid=(3,0)
[Kernel] gid=(4,0)
[Kernel] gid=(0,1)
[Kernel] gid=(1,1)
[Kernel] gid=(2,1)
[Kernel] gid=(3,1)
[Kernel] gid=(4,1)

▶ Sự khác biệt so với khi chạy bằng cấu trúc 1 Chiều Việc định cấu hình số lượng tổng luồng ở mảng gws[] bằng không gian 1 chiều hay 2 chiều sẽ làm thay đổi hoạt động phía GPU ở điểm nào? Đơn giản là nó chỉ làm thay đổi các giá trị trả về khi móc ID qua hàm get_global_id(). Xét trên phương diện Hàm hạt nhân vẫn bị gọi chính xác 10 lần, thì nó hoàn toàn không có mảy may khác biệt nào so với chương trình opencl-mt-1d cả. Tiện thể, nếu bạn sửa đổi phía Mã hạt nhân của chương trình opencl-mt-1d thành như dưới đây, bạn hoàn toàn có thể biến đổi số ID 1 chiều chắt lọc được từ hàm get_global_id() thành các con số của cấu trúc 2 chiều.

uint x_gid = get_global_id(0) % 5;
uint y_gid = get_global_id(0) / 5;

▶ Lợi ích của việc chỉ định bằng 2 Chiều

Cái lợi ích thu được khi chỉ định số lượng luồng bằng tham số không gian 2 chiều lúc gọi hàm clEnqueueNDRangeKernel(), nói một cách cực đoan, chỉ là giúp bạn lược bỏ đi được mấy cái phép tính chia hay lấy phần dư như đoạn mã đã show ở trên. Tuy nhiên, đối với các loại vi xử lý (Processor) thông thường, chúng rất dị ứng với phép chia vì đây là một xử lý tính toán ngốn cực kỳ nhiều thời gian. Việc bớt đi được số lượng phép chia thực hiện bên trong Mã hạt nhân (Kernel code / カーネル・コード) rất có khả năng sẽ dẫn đến sự cải thiện hiệu năng cho Kernel. Ngoài ra, cái ID thu được từ get_global_id() là một con số duy nhất (Unique) được phát cho toàn bộ các luồng (Thread) chạy trên GPU. Vì nó là số ID không đụng hàng trên phạm vi toàn thể GPU, nên nó mới được gắn cái mác “Global” (Toàn cục).

Tính toán lượng lớn dữ liệu bằng Đa luồng (Multithreading)

chap5-2/opencl-mt-add-uint

● Viết chương trình xử lý tính toán một núi dữ liệu

Chỉ để Hàm hạt nhân in mỗi cái Global ID của luồng ra thì chẳng mang tính thực tiễn tẹo nào. Chúng ta sẽ kết hợp cái chương trình có khả năng nhồi nhét lượng dữ liệu lớn giữa Máy chủ và Hạt nhân (Danh sách 3) với cái chương trình cho phép vắt kiệt GPU bằng xử lý Đa luồng (Danh sách 4), để nhào nặn ra một chương trình có thể cân đẹp xử lý tính toán trên một núi dữ liệu thông qua Đa luồng.

▶ Phía Mã máy chủ (Host code)

Ở đây, ta xét đến việc tính toán $C=A+B$ cho tập hợp $N$ khối dữ liệu. Trước tiên, hãy đập Mã máy chủ đi xây lại như Danh sách 6(a). Do số lượng luồng chỉ định ở mảng gws[] đang được thiết lập là 4 theo không gian 1 chiều, Hàm hạt nhân sẽ được kích hoạt 4 lần trên GPU.

 

Danh sách 6: Xử lý tính toán lượng lớn dữ liệu bằng Đa luồng (Multithreading)

				
					// (a) Phía Mã máy chủ (Host code)
#define N_DATA 4

int main(int argc, char *argv[])
{
  size_t data_bytes;
  cl_uint *ha, *hb, *hc;
  cl_mem mem_obj_a, mem_obj_b, mem_obj_c;

  // Cấp phát vùng đệm bộ nhớ ở phía Host
  data_bytes = N_DATA * sizeof(*ha);
  ha = (cl_uint *)malloc(data_bytes);
  hb = (cl_uint *)malloc(data_bytes);
  hc = (cl_uint *)malloc(data_bytes);

  // Ghi dữ liệu đầu vào vào vùng đệm bộ nhớ phía Host
  for (i = 0; i < N_DATA; i++) {
    ha[i] = i;
    hb[i] = i * 100;
  }

  // Thực thi Kernel
  set_work_size(gws, N_DATA, 0);
  set_work_size(lws, 1, 0);
  ret = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, gws, lws, 0, NULL, NULL);

  // Copy dữ liệu đầu ra từ VRAM về bộ nhớ Host (Máy chủ)
  clEnqueueReadBuffer(queue, mem_obj_c, CL_TRUE, 0, data_bytes, hc, 0, NULL, NULL);

  // In kết quả tính toán trên GPU
  for (i = 0; i < N_DATA; i++) {
    printf("[Host] ha[%d]=%u hb[%d]=%u hc[%d]=%u\n", i, ha[i], i, hb[i], i, hc[i]);
  }
}

// (b) Phía Mã hạt nhân (Kernel code)
__kernel void kernel_func(__global uint *a, __global uint *b, __global uint *c)
{
  uint x_gid = get_global_id(0);
  
  c[x_gid] = a[x_gid] + b[x_gid];
}

▶ Hàm hạt nhân

Về phía Mã hạt nhân (Kernel code / カーネル・コード), như ở Danh sách 6(b), bằng cách tận dụng số ID luồng móc được qua lời gọi get_global_id(0) làm luôn số Chỉ mục (Index number) cho mảng dữ liệu a/b/c trên VRAM, mỗi luồng chạy bên phía GPU sẽ gánh một phần công việc trong phép tính $A+B=C$. Hình ảnh minh họa việc Mã hạt nhân được chẻ ra chạy Đa luồng được trình bày tại Hình 4.

Hình 4: Trạng thái thực hiện đa luồng của hàm bên mã hạt nhân

● Kết quả thực thi

Kết quả thực thi được hiển thị bên dưới.

Bash

$ ./main
[Kernel] x_gid=0 : c[0]=a[0]+b[0], a[0]=0, b[0]=0, c[0]=0
[Kernel] x_gid=1 : c[1]=a[1]+b[1], a[1]=1, b[1]=100, c[1]=101
[Kernel] x_gid=2 : c[2]=a[2]+b[2], a[2]=2, b[2]=200, c[2]=202
[Kernel] x_gid=3 : c[3]=a[3]+b[3], a[3]=3, b[3]=300, c[3]=303
  :
[Host] ha[0]=0 hb[0]=0 hc[0]=0
[Host] ha[1]=1 hb[1]=100 hc[1]=101
[Host] ha[2]=2 hb[2]=200 hc[2]=202
[Host] ha[3]=3 hb[3]=300 hc[3]=303

Tôi đã thêm hàm printf() vào Mã hạt nhân để in ra việc hệ thống đang xào nấu phép tính trên phần tử thứ mấy của cái mảng (Array / 配列).

Hàm hạt nhân đã được réo tên 4 lần, và tại mỗi luồng được khởi chạy trên GPU, chúng đang chia nhau thực thi phép tính $A+B=C$ tương ứng cho từng phần tử của mảng dữ liệu đầu vào. Bên cạnh đó, kết quả tính toán do GPU ói ra đã được trả về đúng cách cho phía CPU Máy chủ.

＊

Dù là trong OpenCL hay CUDA, hệ thống luôn dọn sẵn cơ chế để các luồng (Thread) của GPU có được thông tin Global ID, từ đó mỗi luồng sẽ tự giác gánh vác phần việc được giao phó cho mình. Nhờ vào việc hàng hà sa số các Phần tử xử lý (Processor element / プロセッサ・エレメント, PE) (Chú thích 1) ẩn mình trong GPU nhất trí đồng lòng chia chác công việc, toàn bộ bộ máy GPU đã có thể bào ra khối lượng xử lý tính toán với một tốc độ kinh hồn bạt vía.

Chú thích 1: Ở NVIDIA, PE được gọi bằng cái tên CUDA Core, còn bên AMD thì họ gọi nó là Streaming Processor.

3. Các điều kiện ràng buộc của Hạt nhân (Kernel) OpenCL

Những điều Mã hạt nhân không thể làm

So với các chương trình C thông thường chạy trên phía CPU Máy chủ, Mã hạt nhân (Kernel code / カーネル・コード) phải gánh chịu rất nhiều Điều kiện ràng buộc (Constraints / 制約事項) khắt khe.

  1. Tuyệt đối không thể thò tay truy cập vào bộ nhớ phía Máy chủ (Chú thích 2).

  2. Không thể lợi dụng các chức năng của Linux Kernel đang cắm rễ bên phía Máy chủ (ví dụ như gọi các xử lý I/O Tệp – File I/O).

  3. Bó tay trong việc cấp phát động (Dynamic) vùng nhớ VRAM trên GPU trong lúc chương trình đang chạy bằng cách gọi mấy hàm kiểu như malloc().

  4. Một hàm hạt nhân thì không thể nào réo tên một hàm nằm bên phía Máy chủ.

  5. Từ một hàm hạt nhân cũng đừng hòng gọi ra một hàm hạt nhân khác.

  6. Giá trị trả về (Return value) của một hàm hạt nhân là đồ bỏ đi (Không thể sử dụng).

Đối với cái Điều kiện ràng buộc thứ 5, vẫn có phao cứu sinh là xài hàm nội tuyến (Inline function) để thay thế, nên ít ra bạn cũng tránh được thảm cảnh phải lặp đi lặp lại trò copy-paste một đoạn code giống hệt nhau nhiều lần trong mã nguồn (Source code) của Kernel.

Chú thích 2: Trừ một trường hợp ngoại lệ là SVM (Shared Virtual Memory) được bổ sung ở bản OpenCL 2.0. Thế nhưng ngặt nỗi clvk/Ras-Pi lại chưa hỗ trợ món này.

Những điều Mã hạt nhân có thể làm

  1. Truy cập vào vùng nhớ VRAM đã được cấp phát từ trước ở phía Mã máy chủ (Host code / ホスト・コード).

  2. Dư sức thực hiện mọi thể loại phép toán số học (Arithmetic operation) / phép toán logic (Logical operation).

  3. Khai thác các hàm siêu việt (Transcendental function) và hàm tích hợp sẵn (Built-in function) (Chú thích 3).

  4. Sử dụng mọi câu lệnh điều khiển (Control statement) như if/for/while.

    Nhưng ngặt một nỗi, những thiết bị như GPU xử lý phân nhánh điều kiện (Conditional branching) cực kỳ ốc sên khi vướng phải lệnh như if. Nên nếu được, tốt nhất bạn cứ nhào nặn lại bằng các công thức toán học sao cho ra kết quả y chang dùng lệnh if thì hệ thống sẽ cám ơn bạn rất nhiều.

Chú thích 3: Để tường tận xem có những hàm tích hợp sẵn (Built-in function) nào có thể gọi ra xài được, mời bạn tham khảo tài liệu từ trang 74 trở đi tại đường dẫn https://registry.khronos.org/OpenCL/specs/opencl-2.0-openclc.pdf.

Bảng 1 phơi bày các Toán tử và Hàm có thể xài được trong Mã hạt nhân (Kernel code / カーネル・コード). Như tôi có lảm nhảm ở đoạn trước, GPU rất hay hụt hơi với lệnh if, nên nếu chẳng may nghiệp vụ bắt buộc phải chia trường hợp để xử lý, tôi khuyên bạn nên vắt óc cân nhắc xem liệu có thể lôi cái Toán tử 3 ngôi (Ternary operator) ra xài thay cho if được hay không.

Bảng 1: Các Toán tử và Hàm có thể sử dụng trong Mã hạt nhân (Kernel code)

Chủng loại (Type)Ví dụ (Example)Ghi chú (Remarks)
Bốn phép tính số học+, -, *, /Phép chia (Division) tốc độ rùa bò nên ráng mà né đi
Phép toán Bit (Bitwise)&, `, ^, ~`, v.v.
Dịch Bit (Bit shift)>>, <<–
Phép toán logic (Logical)&&, ` 
Phép toán so sánh==, !=, >, <=–
Toán tử 3 ngôi?–
Các loại hàm siêu việtsin, cos, log, exp, v.v.Tốc độ hàm siêu việt cũng không khả quan mấy

Bảng 2 trải ra các Kiểu dữ liệu (Data type) có thể trưng dụng được trong Mã hạt nhân. Về mặt đặc tả kỹ thuật (Specification) của OpenCL thì mấy cái kiểu dữ liệu quốc dân vẫn xài được bình thường thôi. Chẳng qua là trong Mã hạt nhân, việc nhồi nhét Kiểu dữ liệu vector (Vector type / ベクタ型) vào xử lý là một bước đi mang tính sống còn để tối ưu hóa hiệu năng (Performance / 性能).

Bảng 2: Các kiểu dữ liệu có thể sử dụng trong Mã hạt nhân (Kernel code)

Kiểu (Type)Ví dụ (Example)Ghi chú (Remarks)
Số nguyên (Integer) 8 bitchar, uchar–
Số nguyên 16 bitshort, ushort–
Số nguyên 32 bitint, uint–
Số nguyên 64 bitlong, ulong–
Dấu phẩy động (Floating point) 16 bithalf–
Dấu phẩy động 32 bitfloat–
Dấu phẩy động 64 bitdouble–
Các loại Con trỏ (Pointer)uchar *, float *, v.v.–
Các loại Kiểu dữ liệu vector (Vector type)float4, uchar16, v.v.Bắt buộc phải xài nếu muốn tối ưu hóa hiệu năng

Lấy ví dụ, trong 1 cái biến mang dòng máu float4 bạn có thể nhét khít 4 giá trị float vào đó; hay tương tự, biến loại uchar16 sẽ ôm trọn 16 cục giá trị uchar. Số lượng con số được dồn vào 1 cục biến Vector thế này được giang hồ gọi là Độ dài Vector (Vector length). Dành cho bạn nào chưa biết, OpenCL hỗ trợ nhai các độ dài Vector cỡ 2/4/8/16.

Với ngôn ngữ C bình dân hay xài ở Mã máy chủ (Host code), bạn có thể gộp 4 dữ liệu float vào một biến duy nhất bằng chiêu trò khai báo Cấu trúc dữ liệu (Structure / 構造体). Khổ nỗi là, nếu bạn muốn vác cái biến Cấu trúc ấy ra tính toán cái gì đó, bạn sẽ phải è cổ viết ra từng biểu thức tính toán dài dằng dặc cho rành mạch từng biến thành viên (Member variable) trong cái cấu trúc đó.

Trái ngược với điều đó, ở Mã hạt nhân (Kernel code / カーネル・コード), do bạn có thể lôi thẳng cái biến Kiểu dữ liệu vector (Vector type / ベクタ型) ra để giã thẳng các biểu thức tính toán lên đầu nó, nên chỉ với vài ba dòng code, bạn đã vẽ ra được một đoạn chương trình gọn gàng mướt mắt (Bảng 3).

Bảng 3: Cách thức sử dụng Kiểu dữ liệu vector (Vector type) trong OpenCL

Các đoạn mã C truyền thống viết dông dài ở Mã máy chủ (Host code) có thể được cô đọng lại cực kỳ súc tích trong Mã hạt nhân (Kernel code).

Mã máy chủ (Mã C thông thường)Mã hạt nhân (Mã OpenCL C)

typedef struct point_t {

 

float x, y, z, w;

 

} FLOAT4;

Dòng float4 thuộc dạng kiểu dữ liệu có sẵn (Built-in) nên chả cần tốn công khai báo khi xài

FLOAT4 p1;

 

p1.x = 1.0f;

 

p1.y = 2.0f;

 

p1.z = 3.0f;

 

p1.w = 4.0f;

float4 p1;

 

p1.x = 1.0f;

 

p1.y = 2.0f;

 

p1.z = 3.0f;

 

p1.w = 4.0f;

FLOAT4 p1 =

 

{ 1.0f, 2.0f, 3.0f, 4.0f };

float4 p1 =

 

(float4)(1.f, 2.f, 3.f, 4.f);

FLOAT4 p2;

 

p2.x = p1.x * 2.0f;

 

p2.y = p1.y * 2.0f;

 

p2.z = p1.z * 2.0f;

 

p2.w = p1.w * 2.0f;

float4 p2;

 

p2 = p1 * 2.0f;

copy_fl4(

 

float *dest, float *src)

 

{

 

int i;

 

for (i = 0; i < 4; i++) {

 

dest[i] = src[i];

 

}

 

}

copy_fl4(

 

float4 *dest, float4 *src)

 

{

 

*dest = *src;

 

}

LẬP TRÌNH GPU & NPU - Previous Chương 6: Lập trình GPU (GPU Programming / GPUプログラミング) lần đầu tiên … Hello World bằng OpenCL Next - LẬP TRÌNH GPU & NPU Chương 8: Tối ưu hóa ①… Tinh chỉnh (Tuning / チューニング) độ dài vector xử lý

Leave a Comment

Chia sẻ:

Chương 7: Học thông qua thực hành! Lập trình OpenCL

Or copy link

CONTENTS