| 特性 | array<T, N> | array_view<T, N> |
|---|---|---|
| 數據擁有權 | 擁有數據,構造時深拷貝到加速器 | 不擁有數據,僅包裝宿主端指針 |
| 數據傳輸時機 | 構造/賦值時立即上傳 | 核心函數首次訪問時按需上傳 |
| 同步回宿主 | 隱式(析構/同步調用) | 需顯式 synchronize() 或析構時 |
| 適用場景 | 數據生命週期與加速器綁定,重複使用 | 臨時數據、避免不必要拷貝、多加速器共享 |
extent<N> 表示 N 維空間的大小(如 extent<2>(height, width))。parallel_for_each 的計算域、構造 array/array_view。section())、平鋪(tile())。將計算域劃分為固定大小的平鋪(tile),每個平鋪映射到 GPU 的一個線程組(thread block),組內線程共享本地記憶體(共享內存),大幅減少全局內存訪問。
parallel_for_each(extent<2>(H, W).tile<16,16>(),
[=](tiled_index<16,16> idx) restrict(amp) {
// idx.local: 平鋪內局部索引 (0..15, 0..15)
// idx.global: 全局索引
// idx.tile: 平鋪索引
// tile_static 變量:共享內存,僅組內可見
tile_static float shared[16][16];
// 協作加載數據到共享內存
// __syncthreads() 同步
// 計算...
});
tile_barrier::wait() 替代 __syncthreads()。void MatrixMulTiled(array_viewA, array_view B, array_view C) { const int TS = 16; parallel_for_each(C.extent.tile (), [=](tiled_index idx) restrict(amp) { int row = idx.global[0], col = idx.global[1]; tile_static float As[TS][TS], Bs[TS][TS]; float sum = 0; for (int t = 0; t < A.extent[1]; t += TS) { // 協作加載 A 的一塊、B 的一塊到共享內存 As[idx.local[0]][idx.local[1]] = A(row, t + idx.local[1]); Bs[idx.local[0]][idx.local[1]] = B(t + idx.local[0], col); idx.barrier.wait(); for (int k = 0; k < TS; k++) sum += As[idx.local[0]][k] * Bs[k][idx.local[1]]; idx.barrier.wait(); } C[idx.global] = sum; }); }
<amp.h>,命名空間 concurrencyarray 適合長期駐留 GPU 的數據;array_view 適合零拷貝傳遞;extent.tile + tile_static 是釋放 GPU 算力的關鍵,特別是矩陣運算、圖像處理、科學計算等數據復用率高的場景。