Tăng tốc Tổng hợp BEV trên GPU NVIDIA cho các Ứng dụng AI Vật lý

Một mô hình thiết kế ngày càng phổ biến cho xe tự hành (AVs), robot và các hệ thống AI không gian là nhận thức góc nhìn từ trên cao (BEV). Các mô hình BEV chiếu các đặc trưng hình ảnh đa camera vào một lưới chung từ trên xuống, cung cấp cho các mô-đun nhận thức và lập kế hoạch hạ nguồn một bố cục không gian chung để suy luận về làn đường, phương tiện, người đi bộ và không gian trống.

Một thao tác chính trong quy trình này là BEV pooling, thao tác này thu thập các đặc trưng hình ảnh, gán trọng số cho chúng bằng thông tin độ sâu và giảm phân tán chúng thành các ô lưới BEV. Đối với các nhà phát triển, giá trị thực tiễn của nhận thức BEV là nó chuyển đổi nhiều góc nhìn cụ thể của camera thành một biểu diễn thống nhất về không gian của cảnh vật. Thay vì suy luận riêng biệt trên từng ảnh camera, các mô-đun hạ nguồn có thể hoạt động trên một bản đồ đặc trưng thống nhất từ trên xuống được căn chỉnh với thế giới xung quanh xe hoặc robot. BEV pooling là bước làm cho biểu diễn này có thể sử dụng được trong thời gian thực: nó biến các đặc trưng hình ảnh nhận biết độ sâu thành một tensor BEV nhỏ gọn có thể cung cấp cho các tác vụ phát hiện, chiếm dụng, dự đoán quỹ đạo, lập bản đồ và lập kế hoạch.

Về mặt khái niệm, điều này rất đơn giản. Tuy nhiên, trong quá trình triển khai, việc gộp BEV có thể trở thành một nút thắt cổ chai về độ trễ vì nó kết hợp truy cập bộ nhớ không đều, đọc chỉ mục lặp lại, hành vi scatter-reduce và các hiệu ứng bộ nhớ đệm cụ thể của GPU.

Bài đăng này sử dụng BEVPoolV3 làm nghiên cứu điển hình về việc tối ưu hóa việc gộp BEV và các toán tử nặng về thu thập hoặc phân tán khác cho GPU NVIDIA. Nó trình bày một quy trình làm việc thực tế mà bạn có thể áp dụng cho các khối lượng công việc của mình: phân loại chế độ bộ nhớ, loại bỏ lưu lượng phân tán dư thừa, ánh xạ việc triển khai kernel tới GPU đích và xác thực nút thắt cổ chai đang hoạt động bằng NVIDIA Nsight Compute. Kết quả hiệu suất cho thấy tại sao quy trình làm việc này lại quan trọng: cùng một toán tử gộp BEV có thể yêu cầu các chiến lược tối ưu hóa khác nhau tùy thuộc vào việc tập hợp làm việc bị giới hạn bởi DRAM hay chủ yếu nằm trong L2.

BEVPoolV3 giảm độ trễ gộp BEV trên GPU NVIDIA RTX như thế nào?

Các công trình trước đây đã đạt được những tiến bộ quan trọng. BEVPoolV2, được gọi là V2 trong bài viết này, đã giới thiệu một công thức gộp BEV hướng đến triển khai hiệu quả cho các mô hình kiểu BEVDet. CUDA-BEVFusion bao gồm bevpool_half_pack10_kernel, được gọi ở đây là V2+DO, sử dụng duyệt ngoài theo chiều sâu để loại bỏ phần lớn việc tải chỉ mục ngoài ô lặp lại của V2.

BEVPoolV3 tiếp tục hướng tối ưu hóa này với bốn thay đổi bổ sung: giảm tải độ sâu trùng lặp, bản đồ phân tán INT32 năm mảng, các chỉ mục được tính toán trước loại bỏ phép chia số nguyên thời gian chạy và ghi đầu ra sở hữu khoảng.

Bài đăng này sử dụng BEVPoolV3 làm nghiên cứu điển hình về cách tối ưu hóa việc gộp BEV và các toán tử nặng về thu thập hoặc phân tán khác cho GPU NVIDIA. Bạn sẽ tìm hiểu cách phân loại khối lượng công việc gộp BEV theo chế độ bộ nhớ, xác định lưu lượng phân tán dư thừa, ánh xạ việc triển khai kernel tới GPU đích và xác thực nút thắt cổ chai đang hoạt động bằng Nsight Compute. Kết quả hiệu suất trên hai GPU NVIDIA RTX cho thấy tại sao quy trình làm việc này lại quan trọng: cùng một thuật toán gộp BEV có thể bị giới hạn bởi DRAM trên một GPU và phần lớn nằm trong L2 trên GPU khác, đòi hỏi các lựa chọn tối ưu hóa khác nhau.

Bài đánh giá so sánh hai GPU NVIDIA RTX đại diện cho các chế độ bộ nhớ khác nhau: NVIDIA RTX A6000, một GPU NVIDIA Ampere SM86 với bộ nhớ đệm L2 6 MB và không có ISA FP8 gốc, và NVIDIA RTX PRO 6000 Blackwell Max-Q Workstation Edition, một GPU NVIDIA Blackwell SM120 với bộ nhớ đệm L2 128 MB và hỗ trợ FP8 gốc. Cấu hình chuẩn được sử dụng ở đây được lấy từ các mẫu nuScenes thực tế và chứa khoảng 209K điểm phân tán, 80 kênh đặc trưng và tập làm việc gộp BEV 49 MB. Tập làm việc đó vượt quá bộ nhớ đệm L2 của RTX A6000 nhưng vừa với bộ nhớ đệm L2 của RTX PRO 6000 Blackwell Max-Q, khiến RTX A6000 bị giới hạn bởi DRAM và RTX PRO 6000 Blackwell Max-Q phần lớn nằm trong L2 sau khi điền ban đầu.

Hình 1. Gộp BEV nâng cao các đặc trưng hình ảnh đa camera với thông tin độ sâu và giảm các điểm phân tán thành một biểu diễn từ trên xuống chung để phát hiện, dự đoán độ chiếm dụng và lập kế hoạch

Trong cấu hình chuẩn, đường dẫn plugin NVIDIA TensorRT kiểu V2 mất 274.0 µs trên RTX PRO 6000 Blackwell Max-Q. BEVPoolV3 giảm điều đó xuống còn 17.3 µs ở FP16 và 16.4 µs ở FP8. Trên RTX A6000, đường dẫn BEVPoolV3 được điều chỉnh DRAM đạt 90.0 µs. Ngoài tốc độ tăng lên, bài đăng này cho thấy một quy trình làm việc có thể lặp lại để tối ưu hóa các kernel giảm phân tán: phân loại tập làm việc, loại bỏ lưu lượng bộ nhớ dư thừa, khớp hình dạng khởi chạy với GPU đích và xác thực kết quả bằng Nsight Compute.

Hình 2. Tăng tốc đường dẫn plugin TensorRT chuẩn so với V2 FP16. Trên RTX A6000, V3 FP16 đạt 19.31 lần so với V2. Trên RTX PRO 6000 Blackwell Max-Q, V3 FP16 đạt 15.84 lần so với V2 và V3 FP8 đạt 16.71 lần so với V2

Điều kiện tiên quyết

Bài đăng này thảo luận về hành vi kernel CUDA, tích hợp plugin TensorRT và phân tích hồ sơ GPU trong bối cảnh gộp BEV. Các điều kiện tiên quyết hữu ích bao gồm:

Các khái niệm kernel CUDA như lập lịch warp, atomic, tải toàn cục được vector hóa và hành vi bộ nhớ đệm DRAM/L2/L1
Tích hợp plugin TensorRT, đặc biệt là giao diện IPluginV3
Phân tích hiệu năng Nsight Compute để xác thực hành vi bộ nhớ, mức độ chiếm dụng và các nút thắt cổ chai về phát lệnh
Kernel BEV-pooling trong CUDA-BEVFusion là triển khai tham chiếu ngoài chiều sâu trước

Để biết thêm thông tin nền tảng liên quan, hãy xem Hướng dẫn Lập trình CUDA C++, tài liệu plugin TensorRT, các mẫu TensorRT và Hướng dẫn Phân tích hiệu năng Nsight Compute.

Phân loại chế độ bộ nhớ

Bước đầu tiên là phân loại xem tập làm việc của BEV-pooling có vừa với L2 hay không. Trong cấu hình tiêu chuẩn, các mảng chính có tổng cộng khoảng 49 MB, chủ yếu là dữ liệu đặc trưng và đầu ra. Con số duy nhất đó xác định chế độ bộ nhớ: nó lớn hơn bộ nhớ đệm L2 6 MB của RTX A6000, nhưng nhỏ hơn bộ nhớ đệm L2 128 MB của RTX PRO 6000 Blackwell Max-Q.

Hình 3. Phân loại tập làm việc BEV-pooling tiêu chuẩn theo dung lượng L2. Cấu hình tiêu chuẩn được lấy từ các mẫu nuScenes thực tế có tập làm việc khoảng 49 MB. Điều này vượt quá bộ nhớ đệm L2 6 MB trên RTX A6000, vì vậy kernel đi theo đường dẫn bị giới hạn bởi DRAM. Tập làm việc tương tự vừa với bộ nhớ đệm L2 128 MB trên RTX PRO 6000 Blackwell Max-Q, vì vậy kernel chủ yếu nằm trong L2 sau khi điền ban đầu. Lưu ý rằng sơ đồ này mang tính khái niệm và không được vẽ theo tỷ lệ chính xác.

Quyết định vừa/không vừa này thay đổi mục tiêu tối ưu hóa. Trên RTX A6000, việc thu thập đặc trưng và lưu lượng đầu ra tràn ra ngoài L2, vì vậy đường dẫn L2 nhỏ ưu tiên giảm byte và lưu trữ đầu ra truyền luồng bộ nhớ đệm. Trên RTX PRO 6000 Blackwell Max-Q, tập làm việc tiêu chuẩn vừa với L2, vì vậy đường dẫn L2 lớn chuyển sang hiệu quả lệnh, mức độ chiếm dụng, chỉ mục được tính toán trước, tải vector và chuyên môn hóa FP8.

Loại bỏ lưu lượng phân tán dư thừa

Việc scatter-reduce BEV có thể được tóm tắt như sau:

out[ranks_bev[t], c] += depth[ranks_depth[t]] * feat[ranks_feat[t], c];

BEVPoolV2 lặp qua các ô kênh bên ngoài vòng lặp scatter. Đối với C=80 và một ô 8 kênh, các chỉ số scatter giống nhau được tải 10 lần. Điều này tạo ra khoảng 25,1 MB lưu lượng chỉ mục cho các chỉ mục chỉ cần 2,51 MB khi đọc một lần. Một thứ tự vòng lặp ngoài theo chiều sâu khắc phục hầu hết vấn đề đó bằng cách lặp qua từng khoảng BEV trước và tích lũy tất cả các kênh cho khoảng đó trong một lần chạy.

BEVPoolV3 mở rộng hướng tối ưu hóa ngoài theo chiều sâu được sử dụng trong CUDA-BEVFusion bevpool_half_pack10_kernel, được gọi ở đây là V2+DO. V2+DO là một cơ sở hữu ích vì nó đã loại bỏ các lần tải chỉ mục lặp lại ở ngoài ô trong BEVPoolV2 và chứng minh giá trị của việc duyệt dựa trên khoảng. BEVPoolV3 giữ hướng đó và thêm bốn thay đổi triển khai giúp cải thiện khả năng tương thích và hiệu suất trên các chế độ bộ nhớ GPU: giảm tải sâu trùng lặp trong mỗi khoảng; một bản đồ scatter INT32 năm mảng sử dụng ranks_depth, ranks_feat, ranks_bev, interval_starts và interval_lengths; các chỉ mục rõ ràng được tính toán trước loại bỏ phép chia số nguyên thời gian chạy; và các ghi đầu ra thuộc sở hữu của khoảng tránh các atomic so với đường dẫn kiểu V2.

Hình 4. V2+DO loại bỏ hầu hết lưu lượng chỉ mục dư thừa. V3 FP16 giảm thêm chi phí bản đồ scatter được căn chỉnh và áp lực lệnh. V3 FP8 giảm một nửa byte tính năng và đầu ra, điều này hữu ích nhất khi tập làm việc nằm trong L2

Bản đồ scatter năm mảng đặc biệt quan trọng trên các GPU L2 lớn. Việc đóng gói (ranks_depth, ranks_feat, ranks_bev) vào một mảng int3 tạo ra một bản ghi 12 byte. Bố cục đó không thuận tiện cho các giao dịch bộ nhớ được căn chỉnh và không ánh xạ rõ ràng sang tải LDG.128 16 byte. Các mảng INT32 riêng biệt cho phép các luồng liền kề hợp nhất các tải được căn chỉnh và tránh sự ghép nối trường. Tổng số byte logic có thể trông tương tự, nhưng luồng lệnh sạch hơn nhiều.

Triển khai scatter-reduce thuộc sở hữu của khoảng

Trong sản xuất, BEVPoolV3 sử dụng nhiều kernel chuyên dụng, nhưng ý tưởng triển khai cốt lõi dễ hiểu hơn như một bản phác thảo logic nhỏ. Bản đồ scatter được chuẩn bị trước, mỗi khoảng BEV được gán cho một chủ sở hữu, chủ sở hữu duyệt qua các điểm trong khoảng đó, tích lũy các kênh tính năng liên quan và ghi đầu ra một lần.

Cấu trúc này loại bỏ công việc giải mã vòng lặp trong xuất hiện khi bản đồ scatter được đóng gói thành một bản ghi duy nhất. Thay vì tái tạo các chỉ mục tại thời gian chạy, kernel đọc các mảng rõ ràng như ranks_depth, ranks_feat, ranks_bev, interval_starts và interval_lengths.

// 1. Use five precomputed scatter arrays.
// 2. Read explicit indices directly, with no runtime index division.
// 3. Let one interval owner accumulate the output cell.
// 4. Load each depth value once per scatter point in the owner loop.

for each interval iv in parallel:
    start  = interval_starts[iv]
    length = interval_lengths[iv]
    bev    = ranks_bev[start]

    acc[channel_tile] = 0

    for offset in 0 .. length - 1:
        t        = start + offset
        d        = depth[ranks_depth[t]]
        feat_row = ranks_feat[t]

        for c in channel_tile:
            acc[c] += d * feat[feat_row, c]

    out[bev, channel_tile] = acc

Bản phác thảo mã này nắm bắt cấu trúc BEVPoolV3 phổ biến: bản đồ phân tán là rõ ràng, việc giải mã chỉ mục thời gian chạy đã được loại bỏ, độ sâu được tải trong vòng lặp chủ sở hữu khoảng và mỗi ô đầu ra được ghi một lần sau khi tích lũy cục bộ.

Các nhân sản xuất chuyên biệt hóa cấu trúc này cho chế độ bộ nhớ mục tiêu. Trên các GPU L2 nhỏ như RTX A6000, việc triển khai ưu tiên giảm byte, tích lũy FP16 half2 và lưu trữ đầu ra truyền luồng bộ nhớ đệm để tensor đầu ra không đẩy dữ liệu chỉ mục hữu ích ra khỏi L2. Trên các GPU L2 lớn như RTX PRO 6000 Blackwell Max-Q, việc triển khai trước tiên khớp với một bao bọc khởi chạy có độ chiếm dụng cao, sau đó giảm chi phí lệnh bằng các chỉ mục được tính toán trước, tải chỉ mục được vector hóa và các vòng lặp bên trong chuyên biệt cho FP8 nơi tập làm việc cư trú trong L2.

Bất biến thuật toán vẫn giữ nguyên: sở hữu khoảng, tránh giải mã chỉ mục thời gian chạy, tích lũy cục bộ và ghi một lần. Công việc cụ thể theo kiến trúc thay đổi cách bất biến đó được triển khai, chứ không phải những gì toán tử BEV-pooling tính toán.

Hình 5. Độ trễ TensorRT của RTX PRO 6000 Blackwell Max-Q trên sáu cấu hình gộp BEV. V3 FP8 là nhanh nhất trên mọi cấu hình, với mức tăng lớn hơn ở số lượng kênh rộng hơn

Kết quả độ trễ tuyệt đối trên RTX PRO 6000 Blackwell Max-Q cho thấy cách đường dẫn L2 lớn hoạt động trên các số điểm và độ rộng kênh khác nhau. Mẫu tối ưu hóa tương tự cũng đúng trên đường dẫn bị giới hạn bởi DRAM của RTX A6000 khi được đo dưới dạng tăng tốc so với cơ sở V2 FP16. Trên RTX A6000, đường dẫn V3 FP16 được điều chỉnh cho DRAM đạt được mức tăng tốc từ 11 đến 22 lần so với V2 trên các cấu hình đã thử nghiệm. Trên RTX PRO 6000 Blackwell Max-Q, V3 FP8 đạt mức tăng tốc từ 11 đến 42 lần so với V2, với mức tăng lớn nhất xuất hiện ở số điểm lớn hơn và cấu hình kênh rộng hơn.

Hình 6. Tăng tốc chéo GPU so với V2 FP16. RTX A6000 V3 FP16 đạt 11–22× trên các cấu hình đã thử nghiệm, trong khi RTX PRO 6000 Blackwell Max-Q V3 FP8 đạt 11–42× tùy thuộc vào số điểm và độ rộng kênh

Triển khai và xác thực plugin TensorRT

BEVPoolV3 được phơi bày dưới dạng toán tử TensorRT IPluginV3. Plugin này chấp nhận bản đồ phân tán năm mảng cộng với depth và feat, sau đó điều phối kernel thích hợp cho lớp GPU và dtype. Đường dẫn benchmark đã sử dụng các bản dựng ONNX-to-TensorRT và phát lại CUDA Graph với trtexec.

Để xác thực, hãy so sánh với tham chiếu FP64 hoặc một đường dẫn V2 đáng tin cậy hiện có. Kernel được điều chỉnh DRAM trên RTX A6000 đã vượt qua tất cả các phần tử đầu ra được kiểm tra trên sáu cấu hình với atol=1e-2, với lỗi tối đa quan sát được là 0.0065. Trên RTX PRO 6000 Blackwell Max-Q, V2 và V3 tạo ra các đầu ra giống hệt nhau cho các cấu hình đã kiểm tra, cho thấy các thay đổi tối ưu hóa scatter-map và launch đã bảo toàn hành vi số học của đường dẫn tham chiếu.

Ánh xạ thuật toán lên phần cứng

Bốn thay đổi thuật toán BEVPoolV3 có tính di động, nhưng nhân sản xuất phải khớp với nút thắt cổ chai GPU đang hoạt động. Quyết định chính là liệu tập làm việc BEV-pooling có vừa với L2 hay không.

Trên RTX A6000, tập làm việc chuẩn vượt quá L2, vì vậy nhân bị giới hạn bởi lưu lượng DRAM thu thập ngẫu nhiên. Do đó, đường dẫn FP16 ưu tiên giảm byte và bảo toàn bộ nhớ đệm. Việc tăng TILE_C từ 8 lên 16 cắt giảm các lần duyệt tile C=80 từ 10 xuống 5, giảm chi phí vòng lặp và công việc vô hướng lặp lại. Sử dụng tích lũy __half2 với __hfma2 tránh việc mở rộng và đóng gói FP16 sang FP32 không cần thiết. Các bộ nhớ lưu trữ đầu ra truyền luồng bộ nhớ đệm ngăn tensor đầu ra 12.8 MB đẩy các mảng chỉ mục nhỏ hơn nằm trong L2 bị loại bỏ. Sau những thay đổi này, đường dẫn RTX A6000 đạt 90.0 µs trong cấu hình chuẩn, so với 1.738.0 µs đối với V2 FP16.

Trên RTX PRO 6000 Blackwell Max-Q, tập làm việc chuẩn vừa với L2, vì vậy các yếu tố giới hạn chuyển sang việc phát lệnh, mức độ chiếm dụng và độ trễ phụ thuộc. Nhân sản xuất trước tiên khớp với phong bì khởi chạy kiểu V2+DO có mức độ chiếm dụng cao, sau đó loại bỏ chi phí vòng lặp bên trong bằng bản đồ phân tán mảng năm và các chỉ mục được tính toán trước. Điều này tránh phép chia số nguyên thời gian chạy và giảm áp lực tải bản đồ phân tán. Trong cấu hình chuẩn, V3 FP16 đạt 17.3 µs so với 37.8 µs đối với V2+DO FP16, tăng tốc 2.18x ở cùng dtype.

Đường dẫn FP8 chuyên sâu hơn vào trường hợp L2 lớn. Vì dữ liệu tính năng và đầu ra được phục vụ từ L2, việc giảm dtype của chúng có thể mang lại lợi ích độ trễ thực sự. Đường dẫn FP8 sản xuất sử dụng các điểm vào theo số lượng kênh, đóng gói chỉ mục LDG.64 cho C=80, và tải tính năng rộng hơn cho C=128 và C=256. Các tổ hợp tích cực hơn, chẳng hạn như thêm việc mở rộng vòng lặp trên đường dẫn chỉ mục được đóng gói, đã không kết hợp một cách gọn gàng vì chúng làm tăng áp lực thanh ghi và lưu lượng tràn.

The precision ladder có một đích đến thực tế, và đánh giá NVFP4 của chúng tôi giúp làm rõ chính xác mỗi định dạng tỏa sáng ở đâu: chúng tôi đã thử nghiệm một đường dẫn NVFP4 lưu trữ các đặc trưng camera trong E2M1 với các thang đo microblock E4M3 mỗi 16 phần tử trong khi giữ độ sâu và đầu ra ở FP8, và ngay cả với một triển khai được tối ưu hóa mạnh mẽ có các bộ tích lũy __half2 được đóng gói, các hệ số tỷ lệ–độ sâu hợp nhất, và một LUT nửa độ chính xác, chi phí giải mã khiến nó chạy chậm hơn đáng kể so với đường cơ sở FP8.

Việc phân tích bằng Nsight Compute cho thấy kernel được lưu trữ hoàn toàn trong bộ nhớ đệm L2, với mức sử dụng băng thông DRAM thấp và smsp__issue_active dao động thấp hơn nhiều so với thông lượng đỉnh, trong khi đường ống ALU xử lý nhiều lệnh hơn đáng kể so với đường ống FMA.

Điều này cho thấy chế độ scatter-reduce này đã nắm bắt được các lợi ích về hiệu suất byte có sẵn ở FP8, trong khi việc trích xuất nibble bổ sung trên mỗi phần tử, giải mã giá trị và gập tỷ lệ trên mỗi microblock của NVFP4 giới thiệu công việc vòng lặp bên trong mà đường dẫn FP8 tránh được thông qua một lần chuyển đổi FP8 sang nửa vô hướng. Kết quả là một câu chuyện về việc đặt khối lượng công việc rõ ràng: NVFP4 vẫn là một sự phù hợp cực kỳ mạnh mẽ cho các hình dạng nhân ma trận bị giới hạn bởi tính toán chảy qua Tensor Cores thông qua MMA.kind::nvfp4, trong khi đối với các khối lượng công việc scatter-reduce cư trú ở L2, FP8 là lý tưởng trên thang dtype.

Phân tích tương tự áp dụng ngoài việc gộp BEV. Đối với các embedding thưa, voxelization, biểu đồ tần suất, giảm phân đoạn và các toán tử thu thập hoặc phân tán nặng khác, trước tiên hãy phân loại chế độ bộ nhớ, sau đó sử dụng Nsight Compute để xác định xem giới hạn hoạt động là băng thông, phát lệnh hay độ chiếm dụng.

Bảng 1 tóm tắt độ trễ đường dẫn plugin TensorRT của RTX PRO 6000 Blackwell Max-Q, được báo cáo là độ trễ trung vị 100 lần lặp.

Cấu hình	Chiều C	V2 FP16	V2+DO FP16	V3 FP16	V3 FP8	V3 FP8 / V2
small	80	137.8 µs	31.5 µs	12.7 µs	12.6 µs	10.94x
canonical	80	274.0 µs	37.8 µs	17.3 µs	16.4 µs	16.71x
lớn	80	749.9 µs	48.0 µs	27.3 µs	24.9 µs	30.12x
rất lớn	80	1,675.0 µs	61.9 µs	48.0 µs	39.8 µs	42.09x
wide_c128	128	457.3 µs	54.2 µs	21.4 µs	14.8 µs	30.90x
wide_c256	256	880.9 µs	152.3 µs	33.4 µs	22.0 µs	40.04x

Bảng 1. Độ trễ đường dẫn plugin TensorRT trên RTX PRO 6000 Blackwell Max-Q trên nhiều cấu hình mô hình. Các giá trị báo cáo độ trễ trung vị 100 lần lặp cho mỗi lần gọi suy luận/plugin được kiểm chuẩn bằng micro giây cho V2 FP16, V2+DO FP16, V3 FP16 và V3 FP8. Cột cuối cùng cho thấy tốc độ tăng của V3 FP8 so với V2 FP16, được tính toán từ việc giảm độ trễ

Các cân nhắc cho các nền tảng cấp biên

Phân tích tương tự có thể mở rộng cho các nền tảng NVIDIA cấp biên, bao gồm NVIDIA DRIVE AGX Thor. Trong các thử nghiệm ban đầu hướng tới biên, đường dẫn FP16 BEVPoolV3 hoạt động tốt vì các cải tiến cốt lõi—loại bỏ lưu lượng rải dư thừa, tránh giải mã chỉ mục thời gian chạy và sử dụng ghi sở hữu khoảng—là độc lập với kiến trúc.

Tuy nhiên, tốc độ tăng của FP8 không phải là tự động. Trên các mục tiêu cấp biên, kích thước vấn đề nhỏ hơn, hành vi phân cấp bộ nhớ, áp lực thanh ghi và chi phí chuyển đổi FP8 có thể giới hạn hoặc bù đắp lợi ích băng thông dtype lý thuyết. Điều này làm cho FP8 trở thành một tối ưu hóa cụ thể theo nhân và kiến trúc thay vì một sự thay thế sẵn có được đảm bảo cho FP16.

Bắt đầu với tối ưu hóa tổnh hợp BEV

Để áp dụng quy trình làm việc BEVPoolV3 cho nhận thức BEV hoặc khối lượng công việc nặng về thu thập/phân tán của riêng bạn, hãy bắt đầu bằng cách phân tích hiệu năng của toán tử một cách độc lập. Đo kích thước tensor đầu ra, độ sâu, chỉ số phân tán và tính năng, sau đó so sánh tập làm việc tổng thể với dung lượng bộ nhớ đệm L2 của GPU mục tiêu.

Sử dụng NVIDIA Nsight Compute để xác thực xem nút thắt cổ chai đang hoạt động là băng thông bộ nhớ, phát lệnh, độ chiếm dụng hay độ trễ phụ thuộc. Sau đó, chọn chiến lược tối ưu hóa phù hợp với chế độ bộ nhớ: giảm byte và lưu trữ bảo toàn bộ nhớ đệm cho các khối lượng công việc bị giới hạn bởi DRAM, hoặc độ chiếm dụng, các chỉ số được tính toán trước, tải vector và chuyên môn hóa dtype cho các khối lượng công việc cư trú trong L2.

Cách tiếp cận tương tự cũng áp dụng cho các embedding thưa, voxelization, biểu đồ, giảm phân đoạn và các kernel bị giới hạn bộ nhớ không đều khác. Sử dụng kết quả BEVPoolV3 làm hướng dẫn để phân tích hiệu năng toán tử của riêng bạn, chọn chiến lược triển khai phù hợp cho GPU mục tiêu và xác thực kết quả trước khi triển khai thông qua TensorRT. Đối với các tài nguyên liên quan, hãy xem tài liệu TensorRT, Hướng dẫn lập trình CUDA C++, tài liệu Nsight Compute và Diễn đàn Nhà phát triển NVIDIA.

____
Bài viết liên quan