Tổng quan về toán tử gbmv

Toán tử gbmv (General Banded Matrix-Vector Multiplication) thực hiện phép nhân giữa ma trận dải (banded matrix) và vector. Đây là một trong những toán tử cốt lõi trong thư viện đại số tuyến tính cơ bản BLAS. Toán tử này được tối ưu hóa đặc biệt cho đặc tính lưu trữ thưa của ma trận dải, đồng thời hỗ trợ các phép toán chuyển vị và quy nạp song song trên nhiều nhân.

Sản phẩm hỗ trợ: Atlas A5 Training Series / Atlas A5 Inference Series.

Cấu trúc thư mục dự án

├── gbmv
│   ├── CMakeLists.txt        // Tệp cấu hình biên dịch
│   ├── README.md             // Tài liệu hướng dẫn
│   ├── gbmv_test.cpp         // Bộ kiểm thử độ chính xác
│   ├── gbmv_perf_test.cpp    // Bộ kiểm thử hiệu năng
│   └── gbmv_test_utils.h     // Các hàm tiện ích kiểm thử

Mô tả chi tiết toán tử

Chức năng và biểu thức toán học

Toán tử gbmv tính toán phép nhân ma trận dải với vector theo biểu thức:

vec_out = scale_a * op(mat_banded) * vec_in + scale_b * vec_out

Trong đó: mat_banded là ma trận dải, vec_in và vec_out là các vector, scale_a và scale_b là các vô hướng, op(mat_banded) có thể là chính nó hoặc ma trận chuyển vị.

Ma trận dải được lưu trữ theo định dạng LAPACK, yêu cầu leading_dim_a * cols phần tử, với leading_dim_a >= lower_band + upper_band + 1. Phần tử mat_banded(i,j) được lưu tại vị trí upper_band + i - j + j * leading_dim_a, chỉ hợp lệ khi max(0, j-upper_band) <= i <= min(rows-1, j+lower_band).

Giao diện lập trình (API)

cannBlasStatus_t cannBlasBandedMatVecMul(
    cannBlasContext_t ctx,
    cannBlasOpType_t op_type,
    int64_t rows,
    int64_t cols,
    int64_t lower_band,
    int64_t upper_band,
    const float *scale_a,
    const float *mat_banded,
    int64_t leading_dim_a,
    const float *vec_in,
    int64_t stride_in,
    const float *scale_b,
    float *vec_out,
    int64_t stride_out);

Bảng mô tả tham số

Đặc tả toán tử

Hàm Kernel: banded_matvec_kernel

Luồng thực thi toán tử

Dữ liệu đầu vào được thu thập thành bộ nhớ liên tục thông qua GatherStrided ở phía Host, sau đó sao chép sang Device qua H2D. Kernel sẽ chia nhỏ các đoạn cột của ma trận và dữ liệu vector từ GM vào UB để thực hiện tính toán nhân và cộng. Kết quả được quy nạp song song đa nhân về vector đầu ra bằng cách sử dụng SetAtomicAdd. Cuối cùng, kết quả được ghi ngược về Host theo bước nhảy đầu ra thông qua ScatterStrided.

Gọi Kernel

Sử dụng cú pháp gọi kernel <<<>>> để thực thi hàm xử lý trên thiết bị.

banded_matvec_kernel<<<grid_size, block_size>>>(
    rows, cols, lower_band, upper_band, 
    scale_a_val, mat_banded_dev, leading_dim_a, 
    vec_in_dev, stride_in, 
    scale_b_val, vec_out_dev, stride_out
);

Phạm vi kiểm thử

Biên dịch và thực thi

Thực hiện các bước sau tại thư mục gốc của dự án để biên dịch và chạy toán tử.

Cấu hình biến môi trường

Tùy thuộc vào cách cài đặt bộ công cụ CANN, chọn lệnh cấu hình phù hợp:

• Cài đặt mặc định cho người dùng root:

  source /usr/local/Ascend/cann/set_env.sh

• Cài đặt mặc định cho người dùng không phải root:

  source $HOME/Ascend/cann/set_env.sh

• Cài đặt tại đường dẫn tùy chỉnh custom_install_path:

  source ${custom_install_path}/cann/set_env.sh

Chạy kiểm thử

Sử dụng script build để biên dịch và thực thi:

bash build.sh --ops=gbmv --run

Kết quả thành công sẽ hiển thị:

[PASS] gbmv_test