1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21

=================================================================================================
                                   Ringi 3D 架构工坊 · 核心全景
┌───────────────────────────────────────────────────────────────────────────────────────────────┐
│                                                                                               │
│    [Python / User Layer]                [C10 / ATen Layer]             [Hardware / CUDA Layer]│
│   ┌───────────────────────┐            ┌───────────────────────┐      ┌─────────────────────┐ │
│   │   torch.Tensor (x)    │            │   c10::TensorImpl     │      │ c10::StorageImpl    │ │
│   │   Shape: [2, 3]       │──(句柄包装)─►│   Sizes: [2, 3]       │──┐   │ (1D Continuous HBM) │ │
│   │   Stride: [3, 1]      │            │   Strides: [3, 1]     │  │   │ [0][1][2][3][4][5]  │ │
│   │   Offset: 0           │            │   Offset: 0           │  │   │ 0x7f00... (24B)     │ │
│   └───────────────────────┘            └───────────────────────┘  │   └──────────▲──────────┘ │
│                                                                   │(共享物理底账) │             │
│   ┌───────────────────────┐            ┌───────────────────────┐  │              │             │
│   │   torch.Tensor (x.t)  │            │   c10::TensorImpl     │  │              │             │
│   │   Shape: [3, 2]       │──(零拷贝)──►│   Sizes: [3, 2]       │──┘              │             │
│   │   Stride: [1, 3]      │            │   Strides: [1, 3]     │                 │             │
│   │   Offset: 0 (Non-Cont)│            │   Offset: 0           │─────────────────┘             │
│   └───────────────────────┘            └───────────────────────┘                               │
│                                                                                               │
└───────────────────────────────────────────────────────────────────────────────────────────────┘
=================================================================================================

1
2
3
4
5
6
7

# ❌ 原本的代码写法：
# out = attn_weights.matmul(v).transpose(1, 2)
# out = out.reshape(batch_size, seq_len, num_heads * head_dim)

# ❌ 算法同学为了"消除潜在警告"，改写成了：
out = attn_weights.matmul(v).transpose(1, 2).contiguous()
out = out.view(batch_size, seq_len, num_heads * head_dim)

1
2
3
4
5
6
7
8
9

+---------------------------------------------------------------------------------------+
│  物理内存 / 显存硬件的唯一真理：                                                      │
│  硬件只认线性字节地址（Byte Address）：0x00000000 -> 0xFFFFFFFF                         │
│                                                                                       │
│  物理内存就像一条无限延伸的一维"纸带"或"磁带”：                                      │
│  ┌────────┬────────┬────────┬────────┬────────┬────────┬────────┬────────┐             │
│  │ Byte 0 │ Byte 1 │ Byte 2 │ Byte 3 │ Byte 4 │ Byte 5 │ Byte 6 │ Byte 7 │ ...         │
│  └────────┴────────┴────────┴────────┴────────┴────────┴────────┴────────┘             │
+---------------------------------------------------------------------------------------+

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25

+---------------------------------------------------------------------------------------+
│                      PyTorch Tensor 三层物理架构核心关系                               │
│                                                                                       │
│  [第一层：栈上门面]       at::Tensor                                                   │
│                          ├── 包含: c10::intrusive_ptr<c10::TensorImpl> (智能指针)      │
│                          └── 特点: 仅占 8 字节，按值传递极轻量                         │
│                                      │                                                │
│                                      ▼                                                │
│  [第二层：元数据核心]     c10::TensorImpl (堆上对象，记录高维几何视角)                 │
│                          ├── sizes_ / shape_          : [2, 3] (逻辑形状)             │
│                          ├── strides_                 : [3, 1] (物理步长)             │
│                          ├── storage_offset_          : 0 (物理首地址偏移)            │
│                          ├── data_type_ (TypeMeta)    : Float32 (4 Bytes)             │
│                          ├── device_opt_              : CUDA:0                        │
│                          ├── is_contiguous_           : true (布尔缓存)               │
│                          └── storage_                 : c10::Storage (强引用指针)     │
│                                                              │                        │
│                                                              ▼                        │
│  [第三层：物理数据底账]   c10::StorageImpl (堆上对象，管理真正的连续字节内存块)        │
│                          ├── data_ptr_ (DataPtr)      : void* (0x7f88a000)            │
│                          ├── numel_ / nbytes_         : 24 Bytes (6 * 4B)             │
│                          ├── allocator_               : CUDACachingAllocator*         │
│                          └── resizable_               : false                         │
│                                                                                       │
+---------------------------------------------------------------------------------------+

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13

// 摘自 pytorch/aten/src/ATen/core/TensorBody.h
namespace at {
class TORCH_API Tensor {
 protected:
  c10::intrusive_ptr<TensorImpl, UndefinedTensorImpl> impl_;
 public:
  // 大量转发到 impl_ 的内联方法
  IntArrayRef sizes() const { return impl_->sizes(); }
  IntArrayRef strides() const { return impl_->strides(); }
  int64_t storage_offset() const { return impl_->storage_offset(); }
  void* data_ptr() const { return impl_->data(); }
};
}

1
2
3
4
5
6
7

// 摘自 pytorch/c10/core/StorageImpl.h 简化版
struct C10_API StorageImpl : public c10::intrusive_ptr_target {
  DataPtr data_ptr_;       // 封装了 void* 裸指针、自定义释放器 Deleter 以及 Device
  size_t numel_;           // 容量大小（以字节或特定单位计）
  Allocator* allocator_;   // 内存分配器指针（如 CUDACachingAllocator）
  bool resizable_;         // 是否允许动态扩缩容
};

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27

+---------------------------------------------------------------------------------------+
│                          多 Tensor 共享 Storage 物理映射全景                           │
│                                                                                       │
│    Tensor A (原张量)                   Tensor B (切片)             Tensor C (转置)    │
│  ┌───────────────────────┐          ┌───────────────────┐       ┌───────────────────┐ │
│  │ TensorImpl A          │          │ TensorImpl B      │       │ TensorImpl C      │ │
│  │ Shape:   [3, 4]       │          │ Shape:   [2, 2]   │       │ Shape:   [4, 3]   │ │
│  │ Stride:  [4, 1]       │          │ Stride:  [4, 1]   │       │ Stride:  [1, 4]   │ │
│  │ Offset:  0            │          │ Offset:  5        │       │ Offset:  0        │ │
│  └──────────┬────────────┘          └─────────┬─────────┘       └─────────┬─────────┘ │
│             │                                 │                           │           │
│             │                                 │                           │           │
│             └─────────────────────────┐       │       ┌───────────────────┘           │
│                                       ▼       ▼       ▼                               │
│                              ┌───────────────────────────────────┐                    │
│                              │        c10::StorageImpl           │                    │
│                              │        Ref Count: 3               │                    │
│                              │        nbytes: 48 Bytes           │                    │
│                              │        data_ptr: 0x7f00           │                    │
│                              └────────────────┬──────────────────┘                    │
│                                               │                                       │
│                                               ▼                                       │
│    物理内存条/显存: [0][1][2][3] | [4][5][6][7] | [8][9][10][11]                     │
│                     ─── Row 0 ───   ─── Row 1 ───  ──── Row 2 ────                    │
│                                         ▲  ▲                                          │
│                                         └──┴── Tensor B 观察的物理窗口                │
+---------------------------------------------------------------------------------------+

1
2
3
4
5
6
7
8

+---------------------------------------------------------------------------------------+
│  多维坐标 -> 物理地址投影流水线                                                       │
│                                                                                       │
│   (i0, i1, ..., in-1) ──(点积求和 Σ ik * sk)──► Linear Offset ──(+ storage_offset)    │
│                                                        │                              │
│                                                        ▼                              │
│   Physical Memory Address ◄──(Base + Offset * 4B)── Element Index                     │
+---------------------------------------------------------------------------------------+

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13

对于一个 2x3 的二维矩阵：
    ┌             ┐
    │  A00  A01  A02 │
    │  A10  A11  A12 │
    └             ┘

【行优先 (C-Order / PyTorch 默认)】: 按行连贯铺开
物理内存: [ A00, A01, A02, A10, A11, A12 ]
Strides:  [3, 1]

【列优先 (Fortran-Order / MATLAB)】: 按列连贯铺开
物理内存: [ A00, A10, A01, A11, A02, A12 ]
Strides:  [1, 2]

1
2
3
4
5
6
7
8

+---------------------------------------------------------------------------------------+
│  视图算子的底层操作精髓：                                                              │
│                                                                                       │
│  1. 在堆上新 new 一个 c10::TensorImpl 元数据对象；                                    │
│  2. 修改新 TensorImpl 的 sizes_、strides_、storage_offset_；                           │
│  3. 将新 TensorImpl 的 storage_ 指针指向原本的 StorageImpl（原子引用计数 +1）；       │
│  4. 绝不向 GPU 驱动提交任何数据搬运 Kernel（0 FLOPs，0 Byte 显存拷贝）！                │
+---------------------------------------------------------------------------------------+

1
2
3
4
5
6

// PyTorch 底层 transpose 的伪逻辑（非常简单优雅）：
TensorImpl* new_impl = new TensorImpl(original_impl->storage());
new_impl->sizes_[0] = original_impl->sizes_[1];   // 3
new_impl->sizes_[1] = original_impl->sizes_[0];   // 2
new_impl->strides_[0] = original_impl->strides_[1]; // 1
new_impl->strides_[1] = original_impl->strides_[0]; // 3

1
2
3

x = torch.tensor([1, 2, 3])  # Shape: [3], Strides: [1]
# 将 [3] 广播扩展为 [1000000, 3] (一百万行！)
expanded = x.expand(1000000, 3)

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13

+---------------------------------------------------------------------------------------+
│  torch.expand 的 Stride = 0 魔法                                                      │
│                                                                                       │
│  元数据配置：                                                                         │
│  - Shape:   [1000000, 3]                                                              │
│  - Strides: [0, 1]          <=== 关键！第 0 维的步长直接设为 0！                      │
│  - Offset:  0                                                                         │
│  - Storage: [1, 2, 3]       <=== 物理上依然只有 3 个浮点数！                          │
│                                                                                       │
│  物理地址计算：                                                                       │
│  访问 expanded[999999, 1] (第一百万行第 1 列):                                        │
│  Index = 999999 * 0 + 1 * 1 = 1  ===> 直接命中 Storage[1] = 2！                       │
+---------------------------------------------------------------------------------------+

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15

// PyTorch 底层判断 C-Contiguous 的真实实现
static bool compute_contiguous(IntArrayRef sizes, IntArrayRef strides) {
  int64_t dim = sizes.size();
  int64_t expected_stride = 1;
  // 从最右侧维度向左扫描
  for (int64_t i = dim - 1; i >= 0; --i) {
    if (sizes[i] == 0) return true;
    if (sizes[i] == 1) continue; // 大小为 1 的维度不破坏连续性
    if (strides[i] != expected_stride) {
      return false; // 一旦步长不匹配，立即判定为非连续！
    }
    expected_stride *= sizes[i];
  }
  return true;
}

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11

+---------------------------------------------------------------------------------------+
│  Contiguous 判定实例对比表                                                            │
+---------------------+-------------------+---------------------+-----------------------+
│ 操作                | Shape             | Strides             | is_contiguous()       |
+---------------------+-------------------+---------------------+-----------------------+
│ torch.randn(2, 3)   | [2, 3]            | [3, 1]              | True (标准行优先)      |
│ x.transpose(0, 1)   | [3, 2]            | [1, 3]              | False (步长逆序)      |
│ x[:, ::2] (跨列切片)| [2, 2]            | [3, 2]              | False (内存存在空洞)  |
│ x.expand(5, 2, 3)   | [5, 2, 3]         | [0, 3, 1]           | False (第0维步长为0)  |
│ x.unsqueeze(0)      | [1, 2, 3]         | [6, 3, 1]           | True (size=1 维跳过)  |
+---------------------+-------------------+---------------------+-----------------------+

1
2
3

x = torch.randn(2, 3)
xt = x.t()  # Shape: [3, 2], Strides: [1, 3], is_contiguous: False
xt.view(6)  # 💥 抛出 RuntimeError: view size is not compatible with input tensor's size and stride

 1
 2
 3
 4
 5
 6
 7
 8
 9
10

+---------------------------------------------------------------------------------------------------+
│                    .view() vs .reshape() vs .contiguous() 工业级特性对比表                         │
+------------------+-----------------------+-----------------------+--------------------------------+
│ 特性维度         | .view(*shape)         | .reshape(*shape)      | .contiguous()                  │
+------------------+-----------------------+-----------------------+--------------------------------+
│ 是否保证零拷贝   | 是 (100% 零拷贝)      | 否 (可能发生隐式拷贝) | 视连续性而定 (非连续时深拷贝)  │
│ 对非连续输入的反应| 抛出 RuntimeError 崩溃| 自动深拷贝并重排数据  | 分配新内存进行物理重排数据     │
│ 适用场景         | 性能敏感路径、明确连续| 通用形状变换、灵活编程| 必须恢复行优先连续性的算子前置 │
│ 显存分配风险     | 零风险 (不分配显存)   | 有隐式显存暴涨风险    | 有显存二次分配风险             │
+------------------+-----------------------+-----------------------+--------------------------------+

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11

【模式 A：按行遍历（连续，Stride=1）】
CPU 核心读取 Matrix[i][0] 时发生一次 Cache Miss；
硬件自动将包含该元素在内的 64 字节（16 个 FP32 浮点数）整条 Cache Line 载入 L1 Data Cache！
接下来的 15 次访问（Matrix[i][1] 到 Matrix[i][15]）全部直接命中 L1 Cache（~1ns 时延）！
Cache 命中率高达 93.75%！

【模式 B：按列遍历（非连续，Stride=1024）】
CPU 核心读取 Matrix[0][j] 时载入一条 Cache Line（包含 Matrix[0][j..j+15]）；
但下一次访问的是 Matrix[1][j]，跳跃了 1024 个浮点数（4096 字节，直接跨越了 64 个 Cache Line！）；
刚才载入的 15 个预取浮点数完全没用上，下一次访问再次触发 Cache Miss！
Cache 命中率暴跌至接近 0%！CPU 流水线严重停顿（Stall），耗时激增 10~30 倍！

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14

=================================================================================================
【理想情况：张量是 Contiguous 的（Stride=1）】
- 32 个线程的读取地址是连续的：Thread 0 读 addr, Thread 1 读 addr+4, ..., Thread 31 读 addr+124；
- 总共请求 128 字节连续数据；
- GPU 内存控制器（Memory Controller）发起【1 次 128 字节全局内存事务（Transaction）】！
- 物理总线有效利用率 = 128B / 128B = 100%！

=================================================================================================
【恶劣情况：张量是转置非连续的（例如 Stride=1024）】
- 32 个线程读取的地址离散分布在相隔 4KB 的各个角落；
- GPU 内存控制器必须发起【32 次独立的 32 字节内存事务】（总共向 HBM 请求 1024 字节数据）；
- 实际有用的数据只有 128 字节，其余 896 字节全是废弃流量！
- 物理总线有效利用率 = 128B / 1024B = 12.5%！（有效带宽直接缩水 8 倍！）
=================================================================================================

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20

// 假设处理一个 3D 张量，传入形状 (D0, D1, D2) 和步长 (S0, S1, S2)
__global__ void generic_strided_kernel(float* data, int64_t offset, 
                                       int64_t D0, int64_t D1, int64_t D2,
                                       int64_t S0, int64_t S1, int64_t S2, 
                                       int64_t total_elements) {
    int64_t idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx >= total_elements) return;

    // ⚠️ 极其昂贵的一维线性索引反解多维坐标（包含大量整数除法与取模！）
    int64_t i2 = idx % D2;
    int64_t temp = idx / D2;
    int64_t i1 = temp % D1;
    int64_t i0 = temp / D1;

    // 重新用 Strides 计算物理内存偏移
    int64_t physical_offset = offset + i0 * S0 + i1 * S1 + i2 * S2;

    // 访问物理数据
    data[physical_offset] = data[physical_offset] * 2.0f;
}

1
2
3
4

在 BLAS 规范中：
- 矩阵 A 的物理列跨度称为 LDA；
- 矩阵 B 的物理列跨度称为 LDB；
- 矩阵乘法 API: cublasSgemm(handle, transa, transb, m, n, k, alpha, A, lda, B, ldb, beta, C, ldc)

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16

+-------------------------------------------------------------------------------------------------------------------+
│                              Multi-Head Attention 张量元数据全生命周期追踪表                                       │
+------+---------------------------------+---------------------+-----------------------+---------------+------------+
│ 步骤 │ 操作                            │ 输出 Tensor Shape   │ Strides 步长数组      │ is_contiguous │ 物理显存状态│
+------+---------------------------------+---------------------+-----------------------+---------------+------------+
│ 1    │ 输入 X (QKV Projection)         │ [2, 4096, 4096]     │ [16777216, 4096, 1]   │ True          │ 原始连续   │
│ 2    │ q = Q.view(2, 4096, 32, 128)    │ [2, 4096, 32, 128]  │ [16777216, 4096, 128, 1] True          │ 共享(零拷贝)│
│ 3    │ q = q.transpose(1, 2)           │ [2, 32, 4096, 128]  │ [16777216, 128, 4096, 1] ❌ False    │ 共享(零拷贝)│
│ 4    │ k = K.view(...).transpose(1, 2) │ [2, 32, 4096, 128]  │ [16777216, 128, 4096, 1] ❌ False    │ 共享(零拷贝)│
│ 5    │ scores = q @ k.transpose(-2, -1)│ [2, 32, 4096, 4096] │ [536870912, 16777216, 4096, 1] True   │ GEMM新分配 │
│ 6    │ attn = softmax(scores, dim=-1)  │ [2, 32, 4096, 4096] │ [536870912, 16777216, 4096, 1] True   │ In-place/新│
│ 7    │ out = attn @ v                  │ [2, 32, 4096, 128]  │ [16777216, 524288, 128, 1] True       │ BMM新输出  │
│ 8    │ out = out.transpose(1, 2)       │ [2, 4096, 32, 128]  │ [16777216, 128, 524288, 1] ❌ False  │ 共享(零拷贝)│
│ 9    │ out = out.contiguous()          │ [2, 4096, 32, 128]  │ [16777216, 4096, 128, 1]  True        │ ⚠️物理深拷贝│
│ 10   │ out = out.view(2, 4096, 4096)   │ [2, 4096, 4096]     │ [16777216, 4096, 1]   │ True          │ 共享(零拷贝)│
+------+---------------------------------+---------------------+-----------------------+---------------+------------+

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24

import torch

def inspect_tensor(name: str, t: torch.Tensor):
    print(f"=== {name} ===")
    print(f"  Shape (Sizes)     : {list(t.shape)}")
    print(f"  Strides           : {t.stride()}")
    print(f"  Storage Offset    : {t.storage_offset()}")
    print(f"  Data Pointer      : {hex(t.data_ptr())}")
    print(f"  Storage Base Ptr  : {hex(t.storage().data_ptr())}")
    print(f"  Is Contiguous     : {t.is_contiguous()}")
    print(f"  Storage Numel     : {t.storage().nbytes() // t.element_size()}")
    print()

# 创建一个 3x4 的基础张量
x = torch.arange(12, dtype=torch.float32).reshape(3, 4)
inspect_tensor("Original x", x)

# 做一次转置
xt = x.t()
inspect_tensor("Transposed x.t()", xt)

# 做一次切片
x_slice = x[1:, 1:3]
inspect_tensor("Sliced x[1:, 1:3]", x_slice)

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26

=== Original x ===
  Shape (Sizes)     : [3, 4]
  Strides           : (4, 1)
  Storage Offset    : 0
  Data Pointer      : 0x7f9a1000
  Storage Base Ptr  : 0x7f9a1000
  Is Contiguous     : True
  Storage Numel     : 12

=== Transposed x.t() ===
  Shape (Sizes)     : [4, 3]
  Strides           : (1, 4)           <=== 步长对调！
  Storage Offset    : 0
  Data Pointer      : 0x7f9a1000       <=== 与原张量物理指针完全相同！
  Storage Base Ptr  : 0x7f9a1000
  Is Contiguous     : False            <=== 连续性变为 False！
  Storage Numel     : 12

=== Sliced x[1:, 1:3] ===
  Shape (Sizes)     : [2, 2]
  Strides           : (4, 1)
  Storage Offset    : 5                <=== 偏移量精确为 5 (1*4 + 1*1)！
  Data Pointer      : 0x7f9a1014       <=== 0x7f9a1000 + 5*4B = 0x7f9a1014！
  Storage Base Ptr  : 0x7f9a1000       <=== Storage 基地址完全一致！
  Is Contiguous     : False            <=== 跨步切片非连续！
  Storage Numel     : 12

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26

import torch

if not torch.cuda.is_available():
    print("本实验需要 CUDA GPU 环境运行！")
else:
    torch.cuda.empty_cache()
    torch.cuda.reset_peak_memory_stats()

    # 1. 在 GPU 上分配一个 256MB 的张量 (32M 个 FP64)
    print(f"初始已分配显存: {torch.cuda.memory_allocated() / 1024**2:.2f} MB")
    base = torch.randn(4096, 8192, device="cuda", dtype=torch.float32) # 128 MB
    print(f"创建 base 张量后已分配显存: {torch.cuda.memory_allocated() / 1024**2:.2f} MB")

    # 2. 执行 Transpose 转置
    transposed = base.t()
    print(f"执行 base.t() 转置后显存: {torch.cuda.memory_allocated() / 1024**2:.2f} MB (增量: 0 MB!)")

    # 3. 执行 Expand 广播扩大 1000 倍
    expanded = base.unsqueeze(0).expand(1000, 4096, 8192)
    print(f"执行 expand(1000, ...) 后显存: {torch.cuda.memory_allocated() / 1024**2:.2f} MB (增量依然为 0!)")

    # 4. 证明共享物理内存：原地修改 base，观察 transposed 和 expanded
    base[0, 0] = 999.0
    print(f"修改 base[0, 0] 为 999.0 -> transposed[0, 0]: {transposed[0, 0].item()}, expanded[500, 0, 0]: {expanded[500, 0, 0].item()}")
    assert transposed[0, 0].item() == 999.0 and expanded[500, 0, 0].item() == 999.0
    print("✅ 物理共享验证 100% 通过！")

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32

import torch
import time

def benchmark_reshape_vs_contiguous():
    x = torch.randn(4096, 4096, device="cuda" if torch.cuda.is_available() else "cpu")
    xt = x.t() # 非连续

    # 验证 view 报错
    try:
        xt.view(-1)
    except RuntimeError as e:
        print(f"捕获预期异常: {e}")

    # 压测连续情况下的 view (基线)
    start = time.perf_counter()
    for _ in range(1000):
        _ = x.view(-1)
    if torch.cuda.is_available(): torch.cuda.synchronize()
    contiguous_view_time = (time.perf_counter() - start) * 1000

    # 压测非连续情况下的 reshape (触发隐式深拷贝)
    start = time.perf_counter()
    for _ in range(1000):
        _ = xt.reshape(-1)
    if torch.cuda.is_available(): torch.cuda.synchronize()
    non_contiguous_reshape_time = (time.perf_counter() - start) * 1000

    print(f"连续张量 view 1000 次耗时  : {contiguous_view_time:.3f} ms (纯元数据操作)")
    print(f"非连续张量 reshape 1000 次耗时: {non_contiguous_reshape_time:.3f} ms (触发 1000 次深拷贝)")
    print(f"性能差距倍数: {non_contiguous_reshape_time / contiguous_view_time:.1f}x 倍性能劣化！")

benchmark_reshape_vs_contiguous()

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26

import torch
import time

def benchmark_cpu_strided_access():
    size = 4096
    matrix = torch.ones(size, size, dtype=torch.float32)

    # 模式 1：按行求和（连续内存访问，Stride=1）
    start = time.perf_counter()
    sum_row = 0.0
    for i in range(size):
        sum_row += matrix[i, :].sum().item()
    row_time = (time.perf_counter() - start) * 1000

    # 模式 2：按列求和（非连续跨步访问，Stride=4096，每次跳跃 16KB，严重 Cache Miss）
    start = time.perf_counter()
    sum_col = 0.0
    for j in range(size):
        sum_col += matrix[:, j].sum().item()
    col_time = (time.perf_counter() - start) * 1000

    print(f"CPU 按行连续访问求和耗时: {row_time:.2f} ms")
    print(f"CPU 按列跨步访问求和耗时: {col_time:.2f} ms")
    print(f"Cache Miss 带来的耗时惩罚: {col_time / row_time:.2f}x 倍！")

benchmark_cpu_strided_access()

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27

import torch
import time

if torch.cuda.is_available():
    num_elements = 64 * 1024 * 1024 # 64M 个浮点数 (256 MB)
    a = torch.randn(num_elements, device="cuda", dtype=torch.float32)

    # 连续读取求和
    torch.cuda.synchronize()
    start = time.perf_counter()
    for _ in range(100):
        b = a + 1.0 # 连续合并访问
    torch.cuda.synchronize()
    cont_time = (time.perf_counter() - start) / 100

    # 构造非连续跨步切片 (每隔 4 个元素取一个)
    strided_a = a[::4] # 非连续跨步
    torch.cuda.synchronize()
    start = time.perf_counter()
    for _ in range(100):
        b = strided_a + 1.0 # 非合并访存
    torch.cuda.synchronize()
    strided_time = (time.perf_counter() - start) / 100

    print(f"GPU 连续合并访存耗时: {cont_time * 1000:.3f} ms")
    print(f"GPU 跨步非合并访存耗时: {strided_time * 1000:.3f} ms")
    print(f"跨步访存导致算子耗时膨胀: {strided_time / cont_time:.2f}x 倍！")

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25

#include <torch/extension.h>
#include <iostream>

void print_tensor_internals(torch::Tensor t) {
    // 1. 获取 C10 TensorImpl
    c10::TensorImpl* impl = t.unsafeGetTensorImpl();
    
    // 2. 获取维度与步长
    auto sizes = t.sizes();
    auto strides = t.strides();
    int64_t offset = t.storage_offset();
    
    std::cout << "--- C++ ATen Tensor Internals ---" << std::endl;
    std::cout << "Dim: " << t.dim() << ", Offset: " << offset << std::endl;
    std::cout << "Sizes: " << sizes << ", Strides: " << strides << std::endl;
    std::cout << "Is Contiguous: " << t.is_contiguous() << std::endl;
    
    // 3. 安全获取连续内存裸指针
    if (t.is_contiguous()) {
        float* raw_ptr = t.data_ptr<float>();
        std::cout << "Direct Fast Path: raw pointer = " << raw_ptr << std::endl;
    } else {
        std::cout << "Warning: Tensor is Non-Contiguous! Must handle with Strided indexing." << std::endl;
    }
}

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18

+----+---------------------------------------------------+---------------------------------------------------+
| 序号| ❌ 常见小白错误理解                               | ✅ 大厂 AI Infra 正确理解                         |
+----+---------------------------------------------------+---------------------------------------------------+
| 1  | 以为 Tensor 是多维数组，转置和切片会发生数据搬移  | 物理内存永远是一维连续 Storage，转置和切片只是修改|
|    |                                                   | TensorImpl 元数据中的 sizes 和 strides（零拷贝）。 |
+----+---------------------------------------------------+---------------------------------------------------+
| 2  | 遇到 view 报错无脑加 .contiguous()，以为没有代价  | .contiguous() 在非连续时会触发真实的显存分配与全局|
|    |                                                   | 内存深拷贝，在热点循环中会导致 OOM 与吞吐雪崩。   |
+----+---------------------------------------------------+---------------------------------------------------+
| 3  | 以为 .reshape() 和 .view() 完全等价，可以随便混用 | .view() 保证 100% 零拷贝（失败即报错）；.reshape()|
|    |                                                   | 在无法 view 时会隐式回退到深拷贝，带来性能不确定性|
+----+---------------------------------------------------+---------------------------------------------------+
| 4  | 以为广播（expand）会把数据在显存中真实复制成千上万份| expand 仅将对应维度的 Stride 设为 0，物理显存增加量|
|    |                                                   | 恒等于 0，利用数学乘 0 投影实现极致复用。         |
+----+---------------------------------------------------+---------------------------------------------------+
| 5  | 以为 GPU 算力极强，张量连续不连续对性能影响不大   | 非连续张量会彻底粉碎 GPU Warp 的访存合并契约，触发|
|    |                                                   | 数倍的 HBM 事务数，有效显存带宽可能暴跌 80% 以上。|
+----+---------------------------------------------------+---------------------------------------------------+

1
2
3
4
5
6
7
8
9

=================================================================================================
                              Ringi 5 点核心速记口诀（张量底层篇）
=================================================================================================
1. 物理内存一线穿，高维全靠步长牵；
2. 转置切片改标头，零存零拷莫发愁；
3. 行优自右向左算，跨步若逆连贯断；
4. 视图崩溃莫慌乱，滥用连续显存烂；
5. 合并访存性能巅，除模换乘记心间！
=================================================================================================