Đây là bản dịch tiếng Việt của tệp Markdown:

Chương 11: Hướng nghiên cứu GPU về Scalarization và Thực thi Affine

Như được mô tả trong Chương 2, các API tính toán GPU, như CUDA và OpenCL, có một mô hình lập trình giống MIMD cho phép lập trình viên khởi chạy một mảng lớn các luồng scalar lên GPU. Mặc dù mỗi luồng scalar này có thể theo một đường dẫn thực thi duy nhất và có thể truy cập vào các vị trí bộ nhớ tùy ý, trong trường hợp thông thường, chúng đều tuân theo một tập hợp nhỏ các đường dẫn thực thi và thực hiện các hoạt động tương tự.

Sự hội tụ của luồng điều khiển trên các luồng GPU được khai thác trên hầu hết, nếu không phải tất cả, các GPU hiện đại thông qua mô hình thực thi SIMT, nơi các luồng scalar được nhóm lại thành các warp chạy trên phần cứng SIMD (xem Mục 3.1.1). Chương này tóm tắt một loạt các nghiên cứu tiếp tục khai thác sự tương tự của các luồng scalar này thông qua scalarization và thực thi affine.

Nhận thức chính của nghiên cứu này nằm ở việc quan sát cấu trúc giá trị [Kim et al., 2013] giữa các luồng đang thực thi cùng một hạt nhân tính toán. Hai loại cấu trúc giá trị, đồng nhất và affine, được minh họa trong hạt nhân tính toán trong Ví dụ 11.1.

Biến đồng nhất Một biến có cùng một giá trị không đổi cho mỗi luồng trong hạt nhân tính toán. Trong Thuật toán 11.1, biến a, cũng như các hằng số THRESHOLD và Y_MAX_VALUE, đều có giá trị đồng nhất trên tất cả các luồng trong hạt nhân tính toán. Một biến đồng nhất có thể được lưu trữ trong một thanh ghi scalar đơn và được tái sử dụng bởi tất cả các luồng trong hạt nhân tính toán.

Biến affine Một biến có các giá trị là một hàm tuyến tính của ID luồng cho mỗi luồng trong hạt nhân tính toán. Trong Thuật toán 11.1, địa chỉ bộ nhớ của biến y[idx] có thể được biểu diễn dưới dạng một phép biến đổi affine của ID luồng threadIdx.x:

&(y[idx]) = &(y[0]) + sizeof(int) * threadIdx.x;

Biểu diễn affine này có thể được lưu trữ dưới dạng một cặp giá trị scalar, một cơ sở và một bước nhảy, điều này gọn gàng hơn nhiều so với việc mở rộng hoàn toàn.

__global__ void vsadd( int y[], int a ) {
    i
```Đây là bản dịch tiếng Việt của tệp Markdown:
 
```c
int idx = threadIdx.x;
y[idx] = y[idx] + a;
if (y[idx] > THRESHOLD)
    y[idx] = Y_MAX_VALUE;
}

Thuật toán 11.1: Ví dụ về các phép toán vô hướng và affine trong một kernel tính toán (từ [Kim và cộng sự, 2013]).

Có nhiều đề xuất nghiên cứu về cách phát hiện và khai thác các biến đồng nhất hoặc affine trong GPU. Phần còn lại của chương này tóm tắt các đề xuất này trong hai khía cạnh này.

Phát hiện các biến đồng nhất hoặc affine

Có hai cách tiếp cận chính để phát hiện sự tồn tại của các biến đồng nhất hoặc affine trong một kernel tính toán GPU: Phát hiện do Trình biên dịch Điều khiển và Phát hiện thông qua Phần cứng.

Phát hiện do Trình biên dịch Điều khiển

Một cách để phát hiện sự tồn tại của các biến đồng nhất hoặc affine trong một kernel tính toán GPU là thực hiện thông qua một phân tích biên dịch đặc biệt. Điều này có thể xảy ra vì các mô hình lập trình GPU hiện có, CUDA và OpenCL, đã cung cấp các phương tiện cho lập trình viên để khai báo một biến là hằng số trong suốt kernel tính toán, cũng như cung cấp một biến đặc biệt cho ID của luồng. Trình biên dịch có thể thực hiện một phân tích phụ thuộc điều khiển để phát hiện các biến phụ thuộc hoàn toàn vào các hằng số và ID của luồng, và đánh dấu chúng là đồng nhất/affine. Các phép toán chỉ hoạt động trên các biến đồng nhất/affine sau đó là ứng cử viên cho việc tiêu chuẩn hóa.

AMD GCN [AMD, 2012] dựa vào trình biên dịch để phát hiện các biến đồng nhất và các phép toán vô hướng có thể được lưu trữ và xử lý bởi một bộ xử lý vô hướng chuyên dụng.

Asanovic và cộng sự [2013] giới thiệu một phân tích hội tụ và biến thể kết hợp cho phép trình biên dịch xác định các phép toán trong một kernel tính toán tùy ý đủ điều kiện để tiêu chuẩn hóa và/hoặc biến đổi affine. Các chỉ thị trong các vùng hội tụ của một kernel tính toán có thể được chuyển đổi thành các chỉ thị vô hướng/affine. Tại bất kỳ chuyển tiếp nào từ các vùng phân kỳ đến các vùng hội tụ của một kernel tính toán, trình biên dịch chèn một chỉ thị syncwarp để xử lý các phụ thuộc đăng ký do sự phân kỳ điều khiển giữa hai vùng. Asanovic và cộng sự [2013] đã áp dụng ...Đây là bản dịch tiếng Việt của tệp Markdown được cung cấp:

Phân tích này được sử dụng để tạo ra các phép toán vô hướng cho kiến trúc Temporal-SIMT [Keckler và cộng sự, 2011, Krashinsky, 2011].

Tính toán Affine Decoupled (DAC) [Wang và Lin, 2017] dựa trên phân tích trình biên dịch tương tự để trích xuất các ứng viên vô hướng và affine để được tách ra thành một warp riêng biệt. Wang và Lin [2017] tăng cường quá trình này bằng cách phân tích affine phân kỳ, với mục tiêu là trích xuất các chuỗi lệnh