CUTLASS: Implicit GEMM Convolution

已于 2023-12-05 15:37:37 修改
阅读量3k 收藏 8
点赞数 2
分类专栏： GPU DeepLearning CUTLASS 文章标签： 人工智能 GPU CUDA
于 2023-08-15 19:21:06 首次发布
原文链接：https://github.com/NVIDIA/cutlass/blob/main/media/docs/implicit_gemm_convolution.md
版权
DeepLearning 同时被 3 个专栏收录
75 篇文章 10 订阅
订阅专栏
GPU
20 篇文章 4 订阅
订阅专栏
CUTLASS
8 篇文章 1 订阅
订阅专栏

以下内容翻译自：CUTLASS 中的 Implicit GEMM Convolution

"CUTLASS Implicit GEMM API"

CUTLASS Convolution

Implicit GEMM 是将卷积操作表述为 GEMM (广义矩阵-矩阵积)。卷积接受激活张量并对其应用滑动滤波器以产生输出张量。

Introduction

此版本的 CUTLASS 包含几个与卷积相关的工件。

Implicit GEMM Algorithm

二维卷积可以映射到矩阵乘：组建一个包含激活张量元素的卷积矩阵，然后由滤波张量形成的矩阵乘以该矩阵。该算法的最早形式通过通常称为 im2col 的操作显式构造卷积矩阵。生成的矩阵按照滤波器大小复制每个激活元素，消耗额外的存储容量和内存带宽。

隐式 GEMM 算法是 CUDA 中分块、分层 GEMM 计算的一种变体：当数据从全局内存加载到共享内存时，通过周密地更新指针和谓词，它会动态形成卷积矩阵的分块。
一旦在共享内存中构成卷积矩阵，现有组件计算 warp-level GEMM 累加卷积结果并更新输出张量。

本节介绍用于 Turing Tensor Core 的高效隐式 GEMM 卷积 CUDA 内核的结构。

Mapping Convolution to GEMM

前向卷积层计算输出张量y = conv(x, w)其中 x(NHWC)、w(KRSC) 和 y(NPQK) 是4-D 张量。

该计算可以通过以下分析函数来描述。

y[n, p, q, k] = sum_c(sum_r(sum_s( x[n, f(p, r), g(q, s), c] * w[k, r, s, c] )))

其中函数f和g定义如下。

f(p, r) = p * stride_h + R - r - 1 + pad_h
g(q, s) = q * stride_w + S - s - 1 + pad_w

CUTLASS Utilities 中提供了 host 和 device 参考实现。

这个计算可以映射到矩阵乘积的元素，如下所示。

C = gemm(A, B)

其中

  • A 是包含激活的 N H W × R S C \mathrm{NHW}\times \mathrm{RSC} NHW×RSC 行主矩阵
  • B 是包含滤波器的 R S C × K \mathrm{RSC}\times \mathrm{K} RSC×K 列主矩阵
  • C 是包含输出的 N P Q × K \mathrm{NPQ}\times \mathrm{K} NPQ×K 行主矩阵

根据以下关系，输出矩阵 C i j C_{ij} Cij​ 的每个元素对应于输出张量y[n, p, q, k]中的一个元素。

y[n, p, q, k] = Cij

其中，

i = q + Q * (p + P * n)
j = k

这些关系可以倒推如下。

k = j

n = i / (PQ)
residual = i % (PQ)

p = residual / Q
q = residual % Q

在 CRS 上迭代以累加结果的三重循环嵌套也可以通过以下关系被线性化并映射到内部 GEMM K K K 维度（不要与滤波器张量维度 K K K 混淆）。

gemm_k = s + S * (r + R * c)

逆向为

c = gemm_k / (RS)
residual = gemm_k % (RS)

r = residual / S
s = residual % S

给定这些等式，GEMM 三重循环嵌套可以使用张量索引进行扩充，如下所示。

int GEMM_M = N * P * Q;
int GEMM_N = K;
int GEMM_K = C * R * S;

for (int gemm_i = 0; gemm_i < GEMM_M; ++gemm_i) {
  for (int gemm_j = 0; gemm_j < GEMM_N; ++gemm_j) {

    int n = gemm_i / (PQ);
    int npq_residual = gemm_i % (PQ);

    int p = npq_residual / Q;
    int q = npq_residual % Q;

    Accumulator accum = 0;

    for (int gemm_k = 0; gemm_k < GEMM_K; ++gemm_k) {

      int k = gemm_j;

      int c = gemm_k / (RS);
      int crs_residual = gemm_k % (RS);

      int r = crs_residual / S;
      int s = crs_residual % S;

      int h = f(p, r);
      int w = g(q, s);

      ElementA a = tensor_A.at({n, h, w, c});
      ElementB b = tensor_B.at({k, r, s, c});

      accum += a * b;
    }

    C[gemm_i * K + gemm_j] = accum;
  }
}

CUTLASS GEMM implementation 在图块上显式迭代。因此，可以实现一个图块迭代器来分析地计算这些函数并加载适当的元素。然而，由此产生的模运算将是计算密集型的，并且开销将限制以 Turing Tensor Core 为目标的 GEMM 内核的性能。

以下部分描述了如何在以 Tensor Core 为目标的分层 GEMM 内核结构体中实施高效的实现。

CUTLASS Convolution Implementation

为了获得最佳性能，建议使用以下参数：

  • 所有张量都是128位对齐的 NHWC 张量；
  • 通道数 © 是 32 个元素的倍数；
  • 滤波器计数 (K) 是 32 个元素的倍数。

这可实现128位向量内存访问，从而产生高效的 CUDA 内核。通过在conv::kernel::DefaultConv2dFprop中设置 AlignmentA 和 AlignmentB，即使在张量核上也支持较小的对齐，但性能低于128位对齐张量。

CUTLASS Device-level Convolution Operator

CUTLASS 定义了接受大量模板参数的 CUDA C++ 模板，以通过操作、数据类型、图块配置、数学指令和融合输出操作来特化生成的内核。

在 turing_tensorop_conv2dfprop.cu，卷积操作定义如下：

/// Define an Implicit GEMM convolution forward propagation (fprop) kernel
using Conv2dFpropKernel = typename cutlass::conv::kernel::DefaultConv2dFprop<
  ElementInputA,                                          // data type of element a (mapped to activation for fprop)                         
  LayoutInputA,                                           // layout of element a (mapped to activation for fprop)
  ElementInputB,                                          // data type of element b (mapped to filters for fprop)  
  LayoutInputB,                                           // layout of element b (mapped to filters for fprop)
  ElementC,                                               // data type of element c (mapped to output for fprop)
  LayoutC,                                                // layout of element c (mapped to output for fprop)
  ElementAccumulator,                                     // data type of internal accumulation
  MMAOp,                                                  // opcode class tag
  SmArch,                                                 // target SM architecture
  ThreadblockShape,                                       // shape of threadblock tile
  WarpShape,                                              // shape of warp-level GEMM tile
  InstructionShape,                                       // shape of target math instruction
  EpilogueOp,                                             // epilogue operator 
  SwizzleThreadBlock,                                     // optional function to reorder threadblocks for locality
  NumStages,                                              // number of pipeline stages in threadblock-scoped GEMM
  cutlass::arch::OpMultiplyAddSaturate,                   // math operation on data of element a and b
  cutlass::conv::IteratorAlgorithm::kOptimized            // global memory iterator algorithm  
>::Kernel

该模板旨在通用并涵盖所有可行的配置。该示例指定了以下具体数据类型、布局和图块形状。

/// Define an Implicit GEMM convolution forward propagation (fprop) kernel
using Conv2dFpropKernel = typename cutlass::conv::kernel::DefaultConv2dFprop<
  cutlass::int4b_t,                                    // data type of element a (mapped to activation for fprop)                         
  cutlass::layout::TensorNHWC,                         // layout of element a (mapped to activation for fprop)
  cutlass::int4b_t,                                    // data type of element b (mapped to filters for fprop)  
  cutlass::layout::TensorNHWC,                         // layout of element b (mapped to filters for fprop)
  int32_t,                                             // data type of element c (mapped to output for fprop)
  cutlass::layout::TensorNHWC,                         // layout of element c (mapped to output for fprop)
  int32_t,                                             // data type of internal accumulation
  cutlass::arch::OpClassTensorOp,                      // opcode class tag
  cutlass::arch::Sm75,                                 // target SM architecture
  cutlass::gemm::GemmShape<128, 128, 128>,             // shape of threadblock tile
  cutlass::gemm::GemmShape<64, 64, 128>,               // shape of warp-level GEMM tile
  cutlass::gemm::GemmShape<8, 8, 32>,                  // shape of target math instruction
  cutlass::epilogue::thread::LinearCombinationClamp<
    int32_t,                                           // data type of output matrix
    8,                                                 // The number of elements per vectorized
                                                       // memory access. This becomes the vector width of
                                                       // math instructions in the epilogue too.
    int32_t,                                           // Data type of accumulator
    float>;    ,                                       // epilogue operator 
  SwizzleThreadBlock,                                  // optional function to reorder threadblocks for locality
  2,                                                   // number of pipeline stages in threadblock-scoped GEMM
  cutlass::arch::OpMultiplyAddSaturate,                // math operation on data of element a and b
  cutlass::conv::IteratorAlgorithm::kOptimized         // global memory iterator algorithm  
>::Kernel

即使用4位整数输入和输出 (cutlass::int4b_t) 计算2D 卷积前向传播。使用32位整数（int32_t）进行内部累加，并对单精度浮点（float）的输出进行元素线性组合运算。

线程块和线程束级别形状指的是 gemm_api.md 描述的分层分块 GEMM 计算。较大的图块可以更好地重用通过共享内存加载的数据，但启动的 CTA 较少，并且对于较小的问题规模可能无法完全占用 GPU。较小的图块配置实现了较低的峰值利用率，但对于实际工作负载，可以更好地匹配 GPU 内的 SM 数量。

Launching the convolution

以下代码将隐式 GEMM 操作的参数收集到一个结构体中。

//
// Define arguments for CUTLASS Convolution
//

// mode (kCrossCorrelation or kConvolution)
cutlass::conv::Mode mode = cutlass::conv::Mode::kCrossCorrelation;

// Split K dimension into 1 partitions
int split_k_slices = 1;

cutlass::conv::Conv2dProblemSize problem_size(      
    options.input_size,
    options.filter_size,
    options.padding,
    options.conv_stride,
    options.dilation,
    options.output_size(),
    mode,
    split_k_slices);

typename ImplicitGemm::Arguments arguments{
  problem_size,
  tensor_a.device_ref(),
  tensor_b.device_ref(),
  tensor_c.device_ref(),
  tensor_c.device_ref(),
  {options.alpha, options.beta},
};

mode标志指示是计算互相关还是卷积。参数input_size、filter_size、padding、conv_stride和dilation指定输入和输出张量的纬度，并表征问题大小。

参数tensor_a.device_ref()、tensor_b.device_ref()和tensor_c.device_ref()是
CUTLASS TensorRef<>对象，其中包含 GPU 设备内存中张量数据的指针和步幅值。

以下代码初始化并启动设备上的隐式 GEMM 操作。以下代码初始化并在设备上启动隐式GEMM操作。 初始化参数结构体，接着将其用于查询设备端工作区要求，并在需要时在设备内存中分配空间。

然后，使用参数结构体和设备内存中的工作空间初始化隐式 GEMM 对象。该初始化步骤预先计算卷积核使用的内部查找表，必要时也可以清除设备端工作空间。

最后，调用初始化的隐式 GEMM 对象，在设备上启动内核。tensor_c现在包含了隐式 GEMM 的结果。

ImplicitGemm implicit_gemm_op;

// Query workspace size
size_t workspace_size = implicit_gemm_op.get_workspace_size(arguments);

// Allocate workspace memory
cutlass::device_memory::allocation<uint8_t> workspace(workspace_size);

// Initialize the Implicit GEMM object
cutlass::Status status = implicit_gemm_op.initialize(arguments, workspace.get());

if (status != cutlass::Status::kSuccess) {
  /* error */
}

//
// Launch initialized CUTLASS kernel
//

status = implicit_gemm_op();

if (status != cutlass::Status::kSuccess) {
  /* error */
}

该示例演示了如何使用 CUTLASS Utilities 中定义的cutlass::HostTensor<>将输入和输出张量写入 CSV 文件。

  std::ofstream output_workspace(ss.str());

  output_workspace 
    << "Input = \n" << tensor_a.host_view() << "\n\n"
    << "Filters = \n" << tensor_b.host_view() << "\n\n";

  // Copy device memory to host backing store
  tensor_c.sync_host();

  output_workspace << "Computed = \n" << tensor_c.host_view() << std::endl;

CUTLASS Components

CUTLASS 定义了以下 CUDA C++模板来实现隐式 GEMM 卷积，这些模板将在后续章节中进行更详细的描述。
激活图块迭代器将激活图块加载到寄存器中。提供了两种实现方案：

滤波器图块迭代器将滤波器加载到寄存器中。 同样，提供了两种实现方式：

优化迭代器涵盖的改进包括：

  • 在主机上预先计算内核不变的指针增量；
  • 在设备端迭代器上计算 cta 不变掩码谓词；
  • 使用 fast divmod 将 GEMM 维度映射到卷积张量。

例如，一个优化的激活迭代器使用 fast divmod 将 GEMM 的 M 维映射到 NPQ。

流水线主循环将线程块范围的图块从全局内存加载到共享内存中，然后应用 CUTRASS 线程束级 GEMM 操作从共享内存加载并向 Turing Tensor Core 发出指令。

存储到共享内存以及使用 Turing Tensor Core 执行 warp 宽度矩阵乘法运算的操作直接从 CUTLASS GEMM 组件应用。这包括以下几个部分：

Loading Activations and Filters

Implicit GEMM Convolution 算法将 GEMM 的 K 维（CRS 的范围）划分为线程块图块，并将每个线程块图块分配给一个滤波器位置和一个通道区间。迭代完所有滤波器位置后，卷积算法前进到下一个通道区间并从滤波器r=0, s=0开始。

主循环的每次迭代都会计算一个线程块图块的矩阵乘积，如 CUTLASS GEMM implementation 中所述。总而言之，激活和滤波器的线程块图块从全局内存中的张量加载并存储到共享内存中。线程块内的每个线程加载一个或多个向量并共同涵盖整个图块。

下图展示了 Implicit GEMM 主循环的一次特定迭代。线程块中的每个线程都映射到 Activation 和 Filter 张量中的若干元素向量。GEMM 的 M 维度中，每个索引对应于输出张量的唯一(N,P,Q)索引，并且可以基于该索引以及滤波器位置(r,s)来计算指针。

"Convolution Forward Propagation on INT4 data."

体现此功能的 CUTLASS 组件是 Conv2dFpropFilterTileAccessIteratorAnalytic。其构造函数计算 GEMM 的 M 维到 (N, P, Q) 的映射，at()方法将线程执行的每个内存访问的线性偏移映射到的 Activation 张量中。此外，valid()方法计算每个滤波器位置和每个内存访问的有效性，以指示内存访问是在张量的边界内还是在边界外。

operator++() 在连续维度和跨步维度上对线程执行的内存访问进行迭代。

// cutlass/conv/threadblock/conv2d_fprop_activation_tile_access_iterator_analytic.h

// Update iterator to thread's next contiguous, strided memory access
Conv2dFpropActivationTileAccessIteratorAnalytic &operator++() {
  ++iteration_contiguous_;
  if (iteration_contiguous_ < ThreadMap::Iterations::kContiguous) {
    return *this;
  }
  iteration_contiguous_ = 0;
  
  ++iteration_strided_;
  if (iteration_strided_ < ThreadMap::Iterations::kStrided) {
    return *this;
  }
  iteration_strided_ = 0;
 
  return *this;
}

在访问完当前线程块图块的所有入口后，advance()更新指向下一个图块的指针。添加到每个指针的偏移量遵循滤波器位置的遍历，执行以下操作之一：

  • 从滤波器位置(r, s, c)前进到滤波器位置(r, s+1, c)
  • 从滤波器位置(r, S-1, c)前进到滤波器位置(r+1, 0, c)
  • 从滤波器位置(R-1, S-1, c)前进到滤波器位置(0, 0, c+32)

方法advance()主体中的逻辑计算激活 GEMM-A 图块的上述三个更新：

// cutlass/conv/threadblock/conv2d_fprop_activation_tile_access_iterator_analytic.h

// Advance to the next access
void advance() {
  // moves to the next tile
  ++filter_s_;
  if (filter_s_ < problem_size_.S) {
    return;
  }
  filter_s_ = 0;
  
  ++filter_r_;
  if (filter_r_ < problem_size_.R) {
    return;
  }
  filter_r_ = 0;
  
  filter_c_ += Shape::kRow * problem_size_.split_k_slices;
}

类似的逻辑也适用于 Conv2dFpropFilterTileAccessIteratorAnalytic。

为了减少主循环体中的计算开销，可以在主机代码中预先计算指针偏移量，并作为其Params结构体中的查找表提供给 CUDA 内核。如 Conv2dFpropFilterTileAccessIteratorOptimized 所示，将从滤波器位置计算偏移的逻辑提取到了Params构造函数中。

// cutlass/conv/threadblock/conv2d_params.h
struct Conv2dFpropActivationIteratorOptimizedParams<layout::TensorNHWC> {
 ...
// next S
inc_next[0] = conv_sign * (int64_t(layout.stride()[0]) * problem_size.dilation_w) * element_size_bits / 8;

// next R
inc_next[1] = conv_sign * (
    int64_t(layout.stride()[1]) * problem_size.dilation_h
    - (problem_size.S - 1) * layout.stride()[0] * problem_size.dilation_w
  ) * element_size_bits / 8;

// next C
inc_next[2] = (
    threadblock_shape.column() * problem_size.split_k_slices
    - conv_sign * int64_t(problem_size.R - 1) * layout.stride()[1] * problem_size.dilation_h
    - conv_sign * int64_t(problem_size.S - 1) * layout.stride()[0] * problem_size.dilation_w
  ) * element_size_bits / 8;

 ...
}

这允许Conv2dFpropActivationTileAccessIteratorOptimized::advance()中的设备代码中仅执行增量表的简单查找。

// cutlass/conv/threadblock/conv2d_fprop_activation_tile_access_iterator_optimized.h
CUTLASS_HOST_DEVICE
void advance() { 

  int next_idx = 0;
 
  // moves to the next tile
  ++filter_s_;
  if (filter_s_ == problem_size_.S) {
    filter_s_ = 0;
    ++filter_r_;
 
    if (filter_r_ < problem_size_.R) {
      next_idx = 1;
    }
    else {
      filter_r_ = 0;
      next_idx = 2;
    }
  }
  
  add_byte_offset_(params_.inc_next[next_idx]); // in addition to Conv2dFpropActivationTileAccessIteratorAnalytic::advance()

  if (next_idx == 2) {  
    filter_c_ += params_.filter_c_delta;
  }
}

Making use of Tensor Cores

Turing Tensor Core 通过在线程束内的所有线程之间共享数据来高效地计算矩阵乘法累加运算。支持以下操作：

ShapeABC
8x8x32int4b_tint4b_tint32_t
8x8x16int8b_tint8b_tint32_t
16x8x8halfhalfhalf
16x8x8halfhalffloat

从功能上讲，Turing 8x8x32矩阵乘法运算将 A、B 和 C 矩阵分布在线程束内的32个线程上，如下图所示。
"Turing Tensor Op"

CUDA 程序员可以通过 PTX 指令 mma.sync 访问此 Tensor Core 操作。CUTLASS 使用 cutlass/arch/mma_sm75.h 中定义的设备端内在函数封装内联 PTX，如下例所示。

unsigned A;   // eight packed 4-bit integer elements
unsigned B;   // eight packed 4-bit integer elements

int C[2];     // two 32-bit integer elements
int D[2];     // two 32-bit integer elements

asm volatile(
  "mma.sync.aligned.m8n8k32.row.col.s32.s4.s4.s32 {%0,%1}, {%2}, {%3}, {%4,%5};\n"
  : "=r"(D[0]), "=r"(D[1])
  : "r"(A), "r"(B), "r"(C[0]), "r"(C[1]));

为了高效地将数据从共享内存加载到寄存器中，并在线程束之间的分布与上述匹配，Turing GPU 架构引入了 ldmatrix。ldmatrix 是最终的线程束协作指令，因为所有线程都向长度为128位的最多32个行向量提供地址。这些行从共享内存中取出，然后分布在每行四个线程的组中。线程内 SMEM 指针和目标寄存器的排列如下所示。图中突出显示了线程0以强调映射。

"Turing ldmatrix PTX instruction"

Turing Tensor Core 对 INT4数据进行乘法累加运算的计算矩阵的大小为8×8×32个元素。ldmatrix 每次操作最多取32行（或列）。可以发出16个 Tensor Core 操作来实现32×32×32矩阵乘积，并完美地消耗由两个 ldmatrix 指令加载的所有数据，如下图所示。通过增加内存指令数量并发出更多 Tensor Core 操作，可以实现更大的图块，最大可达64x64x32的线程束级矩阵操作。 限制是保存累加器元素的寄存器数量。

"Turing ldmatrix PTX instruction feeding Tensor Core operations"

Shared Memory Layouts

在前两节中，我们描述了如何从全局内存中的激活和滤波器张量加载数据来计算卷积，并且描述了 ldmatrix 和 mma.sync 的组合，以从共享内存中获取数据并发出 Tensor Core 操作。

为了确保数据移动效率，必须注意避免 bank 冲突。CUTLASS 使用置换后的共享内存布局来避免存储到共享内存时发生存 bank 冲突，并使用 ldmatrix 从共享内存高效加载。
下图说明了用于从全局内存加载激活和滤波器线程块图块以及共享内存中的置换布局的线程映射。

"Shared Memory layout used for Turing Tensor Cores"

在图中，一个线程束宽度的内存访问以蓝色突出显示，单个线程加载一个128位向量。全局内存中的图块可以对应于激活或滤波器，并假设是通过四个线程加载连续通道进行“条带挖掘”的。

将共享内存可视化为一个"主行"矩阵，其中八列表示8个128位 bank。 如 CUTLASS GTC 2019 slides、recording 中所述，如果每个线程束满足以下条件，则访问共享内存将不会发生冲突：

  • {T0, T1, …, T7} 不访问同一个128位的 bank
  • {T8, T9, …, T15} 不访问同一个128位的 bank
  • {T16, T17, …, T23} 不访问同一个128位的 bank
  • {T24, T25, …, T31} 不访问同一个128位的 bank

为了实现无冲突存储，共享内存布局重新映射条带挖掘排列以转置向量，并对每个线程指针的列索引应用 XOR 运算。 具体来说，

  int store_column = (lane_id % 8) ^ (lane_id / 8);

这种布局上的转换将有助于从共享内存中读取数据切片，以使用 Tensor Core 计算线程束级矩阵乘法。

下图显示了参与 ldmatrix 指令的前16个线程如何逻辑映射到共享内存中矩阵的c=0..31切片。该切片在代码中称为k_group，因为它对应于线程束级矩阵乘法的相同 K 索引。

"Load kgroup=0 from Shared Memory using ldmatrix"

图中下半部分为共享内存中的物理布局，线程根据 XOR 函数进行行列偏移。通过检查，我们可以观察到没有 bank 冲突，因为 T0 … T7 每个访问唯一的存储体，T8 … T15 也是如此，后续相同。

为了前进到共享内存中的下一个“k-group”，根据以下顺序使用 XOR 操作更新指针：

  • ^1 从k=0前进到k=1
  • ^3 从k=1前进到k=2
  • ^1 从k=2前进到k=3
  • ^3 从k=3前进到k=0

其中第一种转变如下所示。

"Advance to kgroup=1 from Shared Memory using ldmatrix"

CUTLASS warp-level GEMM API 定义了用于从置换共享内存加载数据切片并向 Tensor Core 发出操作的模板。

Updating the Output Tensor

主循环终止后，线程束级 GEMM 的累加器图块存储线程束对输出张量的贡献。然而，线程块内的线程之间的数据分布专门用于使用 Tensor Core 的高效矩阵乘法累加运算，并且不利于全局内存的高效合并操作。需要进行数据重排。

Epilogue 是通过共享内存交换累加器元素、加载输出矩阵或张量的切片、应用线性缩放或偏置等元素操作，以及将结果存储到输出张量的组件。

CUTLASS 将其构建为几个组件：

Unit Tests

单元测试在一个独立的 CUDA 内核中验证上述每个组件的功能行为。这就给人们提供了一个便利的环境

  • 检查模板定义
  • 展示在设备代码中使用这些模板的实例化
  • 断言功能正确性

Convolution unit tests

GEMM unit tests

  • Turing Tensor Core 的线程束范围矩阵乘法：gemm_sm75.cu

Epilogue unit tests

Convolution Example

本节介绍所提供的卷积示例，旨在引导读者了解 Implicit GEMM 卷积的 CUTLASS 实现。

Building and Running the Example

示例 09_turing_tensorop_conv2dfprop 计算前向卷积层，其中输入和输出都是 4-b 整数。示例源代码可见于 examples/09_turing_tensorop_conv2dfprop/turing_tensorop_conv2dfprop.cu。

在构建示例之前，首先执行 quickstart.md 中构建任何 CUTLASS 组件的先决条件步骤。计算能力7.5指的是 Turing 架构，这项工作需要 CUDA 10.2 Toolkit 或更高版本，以使用 Turing Tensor Core 的mma PTX instruction。

$ mkdir build && cd build

$ cmake .. -DCUTLASS_NVCC_ARCHS=75

要构建示例，请从构建目录执行make 09_turing_tensorop_conv2dfprop。

$ make 09_turing_tensorop_conv2dfprop

$ ls examples/09_turing_tensorop_conv2dfprop 
examples/09_turing_tensorop_conv2dfprop

此示例提供了一个简单的命令行界面，用于指定4位整数元素 (cutlass::int4b_t) 的4D 张量的范围，将它们初始化为随机值，并计算卷积层的结果。可选地，输入和输出张量可以保存到.csv文件中，并且可以执行 CUTRASS 主机端参考检查以验证正确性。

使用--help运行可以看到完整的用法语句：

$ ./examples/09_turing_tensorop_conv2dfprop/09_turing_tensorop_conv2dfprop --help
09_turing_tensorop_conv2dfprop example

  This example uses Turing's Tensor Core operators on int4 data types to compute
  forward convolution on tensors of layout NHWC.

Options:

  --help               If specified, displays this usage statement.

  --n <int>            Input tensor extent N
  --h <int>            Input tensor extent H
  --w <int>            Input tensor extent W
  --c <int>            Input tensor extent C
  --k <int>            Filter extent K
  --r <int>            Filter extent R
  --s <int>            Filter extent S

  --alpha <float>      Epilogue scalar alpha
  --beta <float>       Epilogue scalar beta

  --ref-check          If set (true), reference check on the host is computed
  --perf-check         If set (true), performance is measured.
  --benchmark          If set (true), performance benchmarking on several layers and batch-size.
  --iterations <int>   Number of profiling iterations to perform.
  --save-workspace     If set, workspace is written to a text file.
  --tag <string>       String to replicate across the first column in the results table



Examples:

$ ./examples/09_turing_tensorop_conv2dfprop/09_turing_tensorop_conv2dfprop  --n=32 --h=224 --w=224 --c=128 --k=256 --r=1 --s=1

$ ./examples/09_turing_tensorop_conv2dfprop/09_turing_tensorop_conv2dfprop  --n=1 --h=224 --w=224 --c=32 --k=32 --r=3 --s=3 --ref-check

请注意，此示例假设所有张量都是128b 对齐且格式为 NHWC。 因此，对于激活、滤波器和输出，维度 C 必须能被32整除。

如果传递了--benchmark选项，则会针对不同的批处理大小对 ResNet50中的几个层进行分析。此示例输出是在使用 CUDA 10.2编译的 NVIDIA RTX 2080上计算的。

build$ ./examples/09_turing_tensorop_conv2dfprop/09_turing_tensorop_conv2dfprop --benchmark

卷积也可以由 CUTLASS Profiler 运行。

Copyright

Copyright © 2017 - 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
SPDX-License-Identifier: BSD-3-Clause

  Redistribution and use in source and binary forms, with or without
  modification, are permitted provided that the following conditions are met:

  1. Redistributions of source code must retain the above copyright notice, this
  list of conditions and the following disclaimer.

  2. Redistributions in binary form must reproduce the above copyright notice,
  this list of conditions and the following disclaimer in the documentation
  and/or other materials provided with the distribution.

  3. Neither the name of the copyright holder nor the names of its
  contributors may be used to endorse or promote products derived from
  this software without specific prior written permission.

  THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
  AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
  IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
  DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
  FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
  DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
  SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
  CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
  OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
  OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
图波列夫
关注
专栏目录
Im2col+GEMM快速卷积
vectorindex的博客
04-02 703
核心思想：将卷积操作（一个个数据窗口相乘求和）变为矩阵乘法，进行快速计算 如何将普通卷积的操作变成有意义的矩阵相乘？？？ 关键点：普通卷积操作：即①对应位置相乘，之后累加。但如果把w矩阵以及x被w圈起来的部分②拉直（展平），就变成了两个向量。 ③然后向量对应位置相乘之后累加，此形式就变成了两个向量点乘，如果把x被圈的部分全部拉直变成向量，④向量组合乘矩阵。把每一个卷积 核也都拉直，凑起来变成矩阵。利用矩阵来进行多个向量的点乘，⑤变成了矩阵乘法。算出来reshape一下就变回去了。 ①②就是Im2.
旷视MegEngine是如何将31*31的大核卷积计算速度提高10倍的
deephub
05-29 2821
2021年提议的Vision Transformer（VIT）已成为计算机视觉深度学习领域的一个有前途的研究主题。随着VIT的研究变得更加深入，一些研究人员受到VIT的大型接收领域的启发，将卷积网络也改造成具有更大的接收场来提高效率。根据有效的接收场（ERF）理论，ERF的大小与内核大小和模型深度的平方根成正比。这意味着通过堆叠层来实现大ERF的有效性不如增加卷积内核大小。因此，研究人员提出了包含大型卷积内核新的CNN结构。该网络可以达到与VIT相同的准确性。“大内核的cnn可以胜过小内核的cnn”这可能是
CUTLASS学习资料汇总 - CUDA矩阵乘法库
最新发布
m0_75126181的博客
09-26 345
CUTLASS (CUDA Templates for Linear Algebra Subroutines) 是NVIDIA开发的一个CUDA C++模板库,用于在NVIDIA GPU上实现高性能的矩阵乘法(GEMM)和相关计算。它提供了灵活的抽象和可复用的软件组件,可以用来构建自定义的高性能CUDA核函数。支持多种精度的混合计算,包括FP16、BF16、TF32、FP32等针对NVIDIA Tensor Core进行了优化提供了分层的软件抽象,可以灵活地组合不同的计算原语。
Implicit GEMM Algorithm
qq_27763933的博客
11-21 149
conv im2col NV 实现原理。
使用cutlass实现多种精度的GEMM，支持cuda core与tensor core
weixin_42764932的博客
03-27 2052
使用cutlass实现多种精度的GEMM，附有完整代码与cmakelist
Low-memory GEMM-based convolution algorithms for deep neural networks 深度神经网络中基于GEMM的低内存卷积算法
知行合一，止于至善
10-20 1237
文章目录1 相关介绍DNN卷积的实现方法contributions2 DNN 卷积3 使用O(K2CHW)O(K^2CHW)O(K2CHW) 空间的patch matrix的卷积3.1 Matrix layouts3.2 Patch-minor layouts3.3 Patch-building Algoritmsim2col−scanim2col−copy−selfim2col−copy−lon...
CUTLASS的GEMM（low-bit 内存搬运过程记录）（简单版）
qq_37592015的博客
03-25 423
记录CUTLASS的low-bit矩阵乘法GEMM计算中，主机内存到设备内存的搬运过程。
cutlass入门: 调用cutlass做通用矩阵乘法Gemm（附代码）
LconLu的博客
05-28 7547
cutlass是CUDA C++模板抽象的集合，用于实现CUDA中所有级别和规模的高性能矩阵乘法（GEMM）和相关计算。相较于cuBLAS和cuDNN，cutlass中包含了更多可重用的模块化软件组件，这使得cutlass相较于前两者更为灵活。 本文将展示如何用cutlass实现最基本的矩阵计算。 cutlass的使用流程与普通kernel大致相同：先在host端分配空间生成数据，再将host端的数据传入device端的buffer中，输入参数调用cutlass模块进行运算，最后将device端的数据传
cutlass：线性代数子例程的CUDA模板
02-05
CUTLASS是CUDA C ++模板抽象的集合，用于在CUDA中的所有级别和规模上实现高性能矩阵乘法（GEMM）。 它包含类似于用于实现cuBLAS的分层分解和数据移动策略。 CUTLASS将这些“活动部件”分解为C ++模板类抽象的可重用...
Cutlass:具有VulkanGLFWglmstb的C ++ 17简单3D图形库
03-20
Cutlass是一个基于C++17的3D图形库，它整合了多个强大的开源库，如Vulkan、GLFW、glm和stb，为开发者提供了一种简洁高效的方式来创建3D图形应用程序。这个库的设计目标是简化3D图形编程，使得程序员能够更专注于创意...
vim-cutlass：插件，添加了与“删除”分开的“剪切”操作
02-01
《vim-cutlass：扩展剪切功能的Vim插件详解》 在程序员的日常工作中，文本编辑器扮演着至关重要的角色。Vim作为一款历史悠久且功能强大的文本编辑器，深受许多开发者的喜爱。然而，Vim的默认操作中，剪切（Cut）和...
【darknet源码解析-06】gemm.h和gemm.c解析
CVer
09-10 2348
本系列为darknet源码解析，本次解析src/gemm.h 与src/gemm.c两个。在上一篇文章中，我们已经详细讲解了输入特征图如何进行转换，那么在本文中，gemm主要完成矢量和矩阵的加速运算，是darknet卷积底层实现的核心，其实也是caffe卷积实现的核心。 gemm.h 的包含的代码如下：主要就是两个函数的gemm，gemm_cpu的定义【gemm_bin暂不分析】，在这里我们先不...
黑科技：用cutlass进行低成本、高性能卷积算子定制开发
旷视开发者技术社区
09-24 1137
图形处理器通用计算(GPGPU)是指利用 GPU 来计算原本由 CPU 处理的通用计算任务。由于现代 GPU 拥有强大的并行处理能力，通用 GPU 在面对矩阵乘法、卷积等大量并行的计算密...
发表博客之：cutlass 仓库代码中 DefaultGemmConfiguration 细致讲解，对cutlass 感兴趣的老乡站住别走！
fgh123的博客
05-06 688
cutlass
卷积算子优化-2.几种卷积算法
qwesh157的博客
10-14 868
这篇文章来学习一下现在卷积常见的几种实现方法。
im2col+gemm实现卷积基础
LuchangLi 的专栏
05-10 3623
im2col+gemm实现卷积基础 NCHW/NHWC到NC1HWC0数据格式图解
cuDNN 的初始设计
yiran103的专栏
04-22 1199
cuDNN V1.0在2014年的发布,并集成到 Caffe、Paddle 等深度学习框架中。论文介绍了 NVIDIA 对于该库的设计和实现。近十年间，NVIDIA 迭代推出了8代架构，cuDNN 也更新到 8.9。硬件上引入了 Tensor Core，软件方面 cuDNN V8 中的 Graph API 相比之前变化较大。然而，作为深度学习领域影响广泛的软件基础设施，一窥 cuDNN 的初始设计仍是有意义的。cuDNN 类似于 cuBLAS，提供了深度学习原语的高效实现。
im2col的理解
luoshao2000的博客
03-14 738
原文：http://blog.csdn.net/eagelangel/article/details/51287005本人的理解：Im2col之后，图像被整理成了二维的N*(h*w*c)，其中N为卷积窗口数，h,w为卷积尺寸，c为图像通道数，k个卷积核被拉成了(h*w*c)*k，这样我们可以把卷积过程当成一个FC全连接，假设全连接网络的前一层节点为m, 后一层节点为n，则m=h*w*c, n=k，...
RepLKNet实战：使用RepLKNet实现对植物幼苗的分类（非官方）（一）
AI浩
05-25 3157
RepLKNet实战摘要论文解读论文的贡献挑战传统认知整体架构安装包1、安装timm2、安装apex数据增强Cutout和Mixup项目结构计算mean和std生成数据集 摘要 本文通过对植物幼苗分类的实际例子来感受一下超大核的魅力所在。这篇文章能让你学到： 通过对论文的解读，了解RepLKNet超大核的设计思想和架构。 如何使用数据增强，包括transforms的增强、CutOut、MixUp、CutMix等增强手段？ 如何调用自定义的模型？ 如何使用混合精度训练？ 如何使用梯度裁剪防
cutlass实现高性能矩阵乘
08-11
Cutlass是一种用于高性能矩阵乘的开源项目。它是由NVIDIA开发的，旨在利用现代GPU架构的强大计算能力来加速矩阵乘运算。 Cutlass的核心思想是将矩阵乘分解为一系列规模更小且计算复杂度较低的子问题。这些子问题可以在GPU上并行计算，从而提高整个矩阵乘的计算速度。 Cutlass实现矩阵乘主要依赖于CUDA的并行计算模型。它利用GPU的多线程性质，将矩阵的不同部分分配给不同的线程块和线程，以实现并行计算。并且，Cutlass还利用GPU的内存层次结构，通过数据分块和逐块加载的方式，最大限度地提高数据访问效率。 除了并行计算和数据分块，Cutlass还通过调整算法参数和优化内存访问模式等方式，细致地优化矩阵乘运算的性能。同时，Cutlass还提供了一组高效的矩阵乘计算模板，可以根据具体应用场景和硬件设备的特性进行灵活配置和扩展。 总之，Cutlass通过并行计算、数据分块、算法参数调整和内存访问优化等方式，实现了高性能的矩阵乘运算。它利用现代GPU架构的强大计算能力，对于需要大规模矩阵乘的应用程序来说，能够显著提高计算效率和性能表现。
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值