FasterTransformer 004 open_attention.h forward

initialize

在这里插入图片描述

forward()

  • https://github1s.com/NVIDIA/FasterTransformer/blob/v1.0/fastertransformer/cuda/open_attention.h#L149-L217

在这里插入图片描述

  • 使用cuBLAS库执行矩阵乘法运算，并对cublasGemmEx（）进行三个单独的调用。这些操作包括将属性核与输入张量相乘，并添加偏差项，从而生成查询、键和值矩阵。
    在这些矩阵相乘之后，该函数使用sqrtf（）函数计算缩放因子，并将查询、键和值矩阵以及该缩放器传递给另一个名为multiHeadAttr_nofuse_kernelLauncher（）的函数。此函数可能会使用额外的计算将多头注意力应用于查询、键和值矩阵，以生成输出矩阵param_.attr_out。
    最后，forward（）函数捕获执行过程中抛出的任何运行时错误，并重新抛出它们。

cublasGemmEx *3

check_cuda_error(cublasGemmEx(param_.cublas_handle, 
        CUBLAS_OP_N, CUBLAS_OP_N, 
        n, m, k, 
        &alpha, 
        param_.attr_kernel_Q, AType_, n, 
        param_.from_tensor, BType_, k, 
        &beta, 
        query_buf_, CType_, n, 
        computeType_, 
        static_cast<cublasGemmAlgo_t>(cublasAlgo_[0])));

      check_cuda_error(cublasGemmEx(param_.cublas_handle, 
        CUBLAS_OP_N, CUBLAS_OP_N,
        n, m, k, 
        &alpha, 
        param_.attr_kernel_K, AType_, n, 
        param_.to_tensor, BType_, k, 
        &beta, 
        key_buf_, CType_, n, 
        computeType_, 
        static_cast<cublasGemmAlgo_t>(cublasAlgo_[0])));

      check_cuda_error(cublasGemmEx(param_.cublas_handle, 
        CUBLAS_OP_N, CUBLAS_OP_N, 
        n, m, k,
        &alpha,
        param_.attr_kernel_V, AType_, n, 
        param_.to_tensor, BType_, k, 
        &beta, 
        value_buf_, CType_, n, 
        computeType_, 
        static_cast<cublasGemmAlgo_t>(cublasAlgo_[0])));

cublasGemmEx

cublasGemmEx is a function from the NVIDIA cuBLAS library that performs a generalized matrix multiplication operation (GEMM) on two matrices A and B, and accumulates the result into a third matrix C.

The “Ex” suffix in the function name indicates that this is an extended version of the basic cublasGemm function, which allows for more advanced features such as data type casting, tensor operations, and tensor cores support.

The function signature for cublasGemmEx is as follows:

cublasStatus_t cublasGemmEx(cublasHandle_t handle,
                            cublasOperation_t transa,
                            cublasOperation_t transb,
                            int m,
                            int n,
                            int k,
                            const void *alpha,
                            const void *A,
                            cudaDataType_t Atype,
                            int lda,
                            const void *B,
                            cudaDataType_t Btype,
                            int ldb,
                            const void *beta,
                            void *C,
                            cudaDataType_t Ctype,
                            int ldc,
                            cudaDataType_t computeType,
                            cublasGemmAlgo_t algo);

Here’s a brief overview of the input parameters:

  • handle: A handle to the cuBLAS library context.
  • transa and transb: Transpose operation to be performed on matrices A and B respectively before the GEMM operation.
  • m, n, and k: The dimensions of matrices A, B, and C, respectively.
  • alpha: Scalar used to scale the product of matrices A and B.
  • A, B, and C: Pointers to the device memory storing the matrices A, B, and C.
  • Atype, Btype, and Ctype: Data types of matrices A, B, and C, respectively.
  • lda, ldb, and ldc: The leading dimensions of matrices A, B, and C, respectively.
  • beta: Scalar used to scale matrix C before accumulation.
  • computeType: The data type used for intermediate computations.
  • algo: The algorithm used for the GEMM operation.

The function returns a cublasStatus_t value indicating whether the operation was successful or if an error occurred.

multiHeadAttr_nofuse_kernelLauncher

declaration in “open_attention.h”

  void multiHeadAttr_nofuse_kernelLauncher(
      cudaStream_t stream,
      cublasHandle_t handle,
      DataType_* Q,
      const DataType_* bias_Q,
      DataType_* K,
      const DataType_* bias_K,
      DataType_* V,
      const DataType_* bias_V,
      const DataType_* attr_mask,
      DataType_* dst,
      const int batch_size,
      const int seq_len,
      const int head_num,
      const int size_per_head,
      const DataType_ scaler);

define in “open_attention.cu”

  • 定义了一个模板函数和两个特化模板。程序编译时会匹配然后编译。
template void OpenMultiHeadAttention<OperationType::FP32>::multiHeadAttr_nofuse_kernelLauncher(
      cudaStream_t stream,
      cublasHandle_t handle,
      float* Q,
      const float* bias_Q,
      float* K,
      const float* bias_K,
      float* V,
      const float* bias_V,
      const float* attr_mask,
      float* dst,
      const int batch_size,
      const int seq_len,
      const int head_num,
      const int size_per_head,
      const float scaler);

template void OpenMultiHeadAttention<OperationType::HALF>::multiHeadAttr_nofuse_kernelLauncher(
      cudaStream_t stream,
      cublasHandle_t handle,
      __half* Q,
      const __half* bias_Q,
      __half* K,
      const __half* bias_K,
      __half* V,
      const __half* bias_V,
      const __half* attr_mask,
      __half* dst,
      const int batch_size,
      const int seq_len,
      const int head_num,
      const int size_per_head,
      const __half scaler);
}//namespace cuda

template<OperationType OpType_>
void OpenMultiHeadAttention<OpType_>::multiHeadAttr_nofuse_kernelLauncher(
      cudaStream_t stream,
      cublasHandle_t cublas_handle,
      DataType_* Q,
      const DataType_* bias_Q,
      DataType_* K,
      const DataType_* bias_K,
      DataType_* V,
      const DataType_* bias_V,
      const DataType_* attr_mask,
      DataType_* dst,
      const int batch_size,
      const int seq_len,
      const int head_num,
      const int size_per_head,
      const DataType_ scaler)
{

    int m = batch_size * seq_len;
    int k = head_num * size_per_head;

    dim3 grid;
    dim3 block;

    if(OpType_ == OperationType::FP32)
    {
      const int word_per_block = 1;
      assert(k <= 1024);
      assert(m / word_per_block * 3 <= 65536);

      dim3 grid(m / word_per_block * 3);
      dim3 block(k);
      add_QKV_bias<DataType_><<<grid, block, 0, stream>>>(Q, bias_Q, K, bias_K, V, bias_V, q_buf_, k_buf_, v_buf_,
          batch_size, seq_len, head_num, size_per_head, word_per_block);
    }
    else
    {
      const int word_per_block = 1;
      grid.x = batch_size * seq_len / word_per_block;
      block.x = head_num * size_per_head * word_per_block / 2;

      add_QKV_bias<DataType_><<<grid, block, 0, stream>>>(Q, bias_Q, K, bias_K, V, bias_V, q_buf_, k_buf_, 
      v_buf_, batch_size, seq_len, head_num, size_per_head / 2, word_per_block);
    }

    DataType_ alpha = (DataType_)1.0f, beta = (DataType_)0.0f;
    
    check_cuda_error(cublasGemmStridedBatchedEx(cublas_handle,
      CUBLAS_OP_T, CUBLAS_OP_N,
      seq_len, seq_len, size_per_head,
      &alpha,
      k_buf_, AType_, size_per_head, seq_len * size_per_head,
      q_buf_, BType_, size_per_head, seq_len * size_per_head,
      &beta,
      qk_buf_, CType_, seq_len, seq_len * seq_len,
      batch_size * head_num,
      computeType_,
      static_cast<cublasGemmAlgo_t>(cublasAlgo_[1])));

    if(seq_len <= 32)
      block.x = 32;
    else if(seq_len > 32 && seq_len <= 64)
      block.x = 64;
    else if(seq_len > 64 && seq_len <= 128)
      block.x = 128;
    else if(seq_len > 128 && seq_len <= 256)
      block.x = 256;
    else if(seq_len > 256 && seq_len <= 512)
      block.x = 512;
    else
      block.x = 1024;

    if(batch_size * head_num <= 120)
    {
      grid.x = batch_size * head_num * seq_len;
      softmax_kernel_v2<DataType_><<<grid, block, 0, stream>>>(qk_buf_, attr_mask, batch_size, head_num, seq_len, scaler); 
    }
    else
    {
      grid.x = batch_size * head_num;
      softmax_kernel<DataType_><<<grid, block, 0, stream>>>(qk_buf_, attr_mask, batch_size, head_num, seq_len, scaler); 
    }

    check_cuda_error(cublasGemmStridedBatchedEx(cublas_handle,
      CUBLAS_OP_N, CUBLAS_OP_N,
      size_per_head, seq_len, seq_len,
      &alpha,
      v_buf_, AType_, size_per_head, seq_len * size_per_head,
      qk_buf_, BType_, seq_len, seq_len * seq_len,
      &beta,
      transpose_dst_, CType_, size_per_head, seq_len * size_per_head,
      batch_size * head_num,
      computeType_,
      static_cast<cublasGemmAlgo_t>(cublasAlgo_[2])));

/* for half2 only */
    if(OpType_ == OperationType::HALF)
    {
      const int seq_per_block = 4;
      grid.x = batch_size * head_num * seq_len / seq_per_block;
      block.x = seq_per_block * size_per_head / 2;

      assert(grid.x * seq_per_block == batch_size * head_num * seq_len);

      transpose<DataType_><<<grid, block, 0, stream>>>(transpose_dst_, dst, 
          batch_size, seq_len, head_num, size_per_head / 2);
    }
    else
    {
      const int seq_per_block = 1;
      grid.x = batch_size * head_num * seq_len / seq_per_block;
      block.x = seq_per_block * size_per_head;
      transpose<DataType_><<<grid, block, 0, stream>>>(transpose_dst_, dst, 
          batch_size, seq_len, head_num, size_per_head);
    }
}

cublasGemmTridedBatchedEx

cublasGemmTridedBatchedEx是cuBLAS库中的一个函数，用于执行跨步分批矩阵乘法。它获取多组输入矩阵，并对每组并行执行相同的矩阵乘法运算，将结果存储回内存。函数名称中的“Ex”表示这是基本“cublasGemmTriedBatched”函数的扩展版本，其中包括用于指定数据类型、比例因子和其他设置的附加选项。
在深度学习的背景下，“cublasGemmTridedBatchedEx”通常用于在Transformer神经网络中执行多头自注意操作。此操作涉及将查询、键和值矩阵集相乘以产生注意力分数，注意力分数用于对值进行加权并计算最终输出。有效地执行这些矩阵乘法对于在大型Transformer模型中实现高性能至关重要。

Attention & Transformer

在这里插入图片描述
在这里插入图片描述

  • FlashAttention：一种具有 IO 感知，且兼具快速、内存高效的新型注意力算法（但非传统attention的近似方法，是数学等价的）
    https://www.bilibili.com/video/BV1SW4y1X7kh/?
    https://zhuanlan.zhihu.com/p/618533434

CG

cuda inline

  • 您可能会注意到的第一件事是__inline__声明。这可能是不必要的。它告诉编译器 将此函数的整个代码放在该点 调用它的位置，而不是导致跳转发生。 这使得编译后的代码运行得更快。另一方面 NVCC知道使许多（也许是大多数）设备功能内联 已经没有问了，所以我们的要求可能是多余的。 我把它放进去，这样我就可以谈论它了。如果函数 更长，也许默认情况下它不会内联，但是 如果您编写该函数主要是为了可读性，则 使其内联可能对您很重要。在这种情况下， 您可以获得更具可读性的代码，而不会牺牲速度来跳转。
  • 在这里插入图片描述

layernorm

  • https://github1s.com/NVIDIA/FasterTransformer/blob/v1.0/fastertransformer/bert_encoder_transformer.h#L208-L281 中调用了add_bias_input_layernorm_kernelLauncher方法。

  • 一个帮助理解multiHeadAttr_nofuse_kernelLauncher的例子

#include <stdio.h>
#include <iostream>
#include <typeinfo>

// using OperationType = int;
enum class OperationType{FP32, HALF};


template<OperationType op>class First{};
template< OperationType N,template<OperationType>class XXX >
class Second;

template< OperationType N,template<OperationType>class XXX >
class Second{
    public:
        XXX<N> b; 
        Second(){ 
            std::cout<<"NNN"; 
        } 
};

// template<template<OperationType> class MultiHeadAttention_>
// class BertEncoderTransformerTraits<OperationType::FP32, MultiHeadAttention_>
template< template<OperationType>class XXX >
class Second<OperationType::FP32,XXX>{
    public:
        XXX<OperationType::FP32> b; 
        Second(){ 
            std::cout<<"SSSSS"; 
        } 
};

// template< template<OperationType>class XXX >
// class Second<OperationType::HALF,XXX>{
//     public:
//         XXX<OperationType::HALF> b; 
//         Second(){ 
//             std::cout<<"HALF"; 
//         } 
// };




int main()
{
    printf("Hello World\n");
    // Second<OperationType::FP32,First> *second = new  Second<OperationType::FP32,First>();
    Second<OperationType::HALF,First> *second = new  Second<OperationType::HALF,First>();
    return 0;
}

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.mfbz.cn/a/28469.html

如若内容造成侵权/违法违规/事实不符，请联系我们进行投诉反馈qq邮箱809451989@qq.com，一经查实，立即删除！

相关文章

【社区图书馆】《看漫画学Python：有趣、有料、好玩、好用（全彩修订版）》

【社区图书馆】《看漫画学Python：有趣、有料、好玩、好用（全彩修订版）》

背景 Python是一门既简单又强大的编程语言&#xff0c;被广泛应用于数据分析、大数据、网络爬虫、自动化运维、科学计算和人工智能等领域。Python也越来越重要&#xff0c;成为国家计算机等级考试科目&#xff0c;某些中小学也开设了Python编程课程。本书秉承有趣、有料、好玩…
阅读更多...
SpringCloud服务注册与发现组件Eureka(五)

SpringCloud服务注册与发现组件Eureka(五)

Eureka github 地址&#xff1a; https://github.com/Netflix/eureka Eureka简介 Eureka是Netflix开发的服务发现框架&#xff0c;本身是一个基于REST的服务&#xff0c;主要用于定位运行在AWS域中的中间层服务&#xff0c;以达到负载均衡和中间层服务故障转移的目的。Spring…
阅读更多...
【Android -- JNI 和 NDK】JNI 基础知识以及如何使用

【Android -- JNI 和 NDK】JNI 基础知识以及如何使用

JNI 基础知识 我们来系统梳理一下JNI中涉及的基本知识。 JNI定义了以下数据类型&#xff0c;这些类型和Java中的数据类型是一致的&#xff1a; Java原始类型&#xff1a;jint, jbyte, jshort, jlong, jfloat, jdouble, jchar, jboolean这些分别对应这 java 的int, byte, shor…
阅读更多...
css 包含块

css 包含块

你不知道的 CSS 之包含块 一说到 CSS 盒模型&#xff0c;这是很多小伙伴耳熟能详的知识&#xff0c;甚至有的小伙伴还能说出 border-box 和 content-box 这两种盒模型的区别。 但是一说到 CSS 包含块&#xff0c;有的小伙伴就懵圈了&#xff0c;什么是包含块&#xff1f;好像…
阅读更多...
微服务springcloud 02 创建项目中的三个service子系统，springcloud中注册中心Eureka介绍和把三个系统注册到Eureka中

微服务springcloud 02 创建项目中的三个service子系统，springcloud中注册中心Eureka介绍和把三个系统注册到Eureka中

item service项目 01.使用springboot创建项目 02.选择依懒项在这里插入代码片 spring web 03.添加sp01-commons依赖 在pom.xml文件中 <?xml version"1.0" encoding"UTF-8"?> <project xmlns"http://maven.apache.org/POM/4.0.0" x…
阅读更多...
【色度学】光学基础

【色度学】光学基础

1. 光的本质 &#xff08;1&#xff09;波长不同的可见光&#xff0c;引起人眼的颜色感觉不同。 &#xff08;2&#xff09;人们观察到的颜色是物体和特有色光相结合的结果&#xff0c;而不是物体产生颜色的结果。 2. 光度量 【ISP】光的能量与颜色&#xff08;1&#xff0…
阅读更多...
NIO 基础

NIO 基础

3. 文件编程 non-blocking io 非阻塞 IO 1.1 Channel & Buffer channel 类似于 stream&#xff0c;它就是读写数据的双向通道&#xff0c;可以从 channel 将数据读入 buffer&#xff0c;也可以将 buffer 的数据写入 channel&#xff0c;而之前的 stream 要么是输入&#…
阅读更多...
统信UOS V20 安装mysql5.7.42详细教程

统信UOS V20 安装mysql5.7.42详细教程

1 安装包准备 到mysql官网可以看到最新的是8.0.33&#xff0c;想下载其他版本的点击 Looking for previous GA versions?Select Operating System: 选择如下版本的mysql 安装包 2 安装 2.1 上传文件至服务器 下载后通过远程将安装包上传至服务器&#xff0c;我这里将安装…
阅读更多...
Seesion会话超时时间测试-业务安全测试实操(3)

Seesion会话超时时间测试-业务安全测试实操(3)

Seesion会话超时时间测试, Cookie仿冒测试, 密文比对认证测试 本地加密传输测试-业务安全测试实操(2)_luozhonghua2000的博客-CSDN博客 测试原理和方法 在用户成功登录系统获得Session认证会话后，该Session认证会话应具有生命周期，即用户在成功登录系统后，如果在固定时间内…
阅读更多...
两个链表相加

两个链表相加

描述 假设链表中每一个节点的值都在 0 - 9 之间&#xff0c;那么链表整体就可以代表一个整数。 给定两个这种链表&#xff0c;请生成代表两个整数相加值的结果链表。 数据范围&#xff1a;0≤n,m≤1000000&#xff0c;链表任意值 0≤val≤9 要求&#xff1a;空间复杂度 O(n)…
阅读更多...
Triton教程 -- 利用Triton部署你自己的模型

Triton教程 -- 利用Triton部署你自己的模型

Triton教程—利用Triton部署你自己的模型 给定一个经过训练的模型&#xff0c;我如何使用 Triton 推理服务器以最佳配置大规模部署它&#xff1f; 本文档旨在帮助回答这个问题。 对于那些喜欢高级概述的人&#xff0c;下面是大多数用例的通用流程。 对于那些希望直接进入的人…
阅读更多...
Windows Server AD域控服务器升级/迁移（AD域控的五大角色转移）

Windows Server AD域控服务器升级/迁移（AD域控的五大角色转移）

Windows Server AD域控服务器升级/迁移&#xff08;AD域控的五大角色转移&#xff09; 新域控服务器安装配置域控服务器&#xff0c;加入现有域域控角色迁移到新域控服务器原域控服务器降级退域 本文主要介绍在现有域环境下如何进行域控服务器的迁移/升级操作。对于域结构的网络…
阅读更多...
抖音seo矩阵系统源码|需求文档编译说明（一）

抖音seo矩阵系统源码|需求文档编译说明（一）

抖音seo矩阵系统文章目录技术囊括 ①产品原型 ②需求文档 ③产品流程图 ④部署方式说明 ⑤完整源码 ⑥源码编译方式说明 ⑦三方框架和SDK使用情况说明和代码位置 ⑧平台操作文档 ⑨程序架构文档 短视频矩阵系统源码开发锦囊囊括前言一、短视频账号矩阵系统开发者必备能力语言&…
阅读更多...
深度相机介绍

深度相机介绍

一、什么是深度相机 &#xff08;五&#xff09;深度相机&#xff1a;结构光、TOF、双目相机 - 知乎 传统的RGB彩色普通相机称为2D相机&#xff0c;只能拍摄相机视角内的物体&#xff0c;没有物体到相机的距离信息&#xff0c;只能凭感觉感知物体的远近&#xff0c;没有明确的数…
阅读更多...
基于SpringBoot+vue的简历系统设计和实现

基于SpringBoot+vue的简历系统设计和实现

博主介绍&#xff1a; 大家好&#xff0c;我是一名在Java圈混迹十余年的程序员&#xff0c;精通Java编程语言&#xff0c;同时也熟练掌握微信小程序、Python和Android等技术&#xff0c;能够为大家提供全方位的技术支持和交流。 我擅长在JavaWeb、SSH、SSM、SpringBoot等框架下…
阅读更多...
seatunnel入门案例，集群模式

seatunnel入门案例，集群模式

目录 安装部署 解压 环境变量 安装plugin 添加资源jar包 SEATUNNEL 配置文件 env&#xff1a;环境设置 source&#xff1a;数据源设置 sink&#xff1a;数据去向设置 transform: 数据转换设置 运行方式 seatunnel 引擎(zeta) 本地模式 集群模式 安装部署 解压 tar…
阅读更多...
深入浅出Node.js中的node_modules

深入浅出Node.js中的node_modules

文章目录 1. 什么是node_modulesnode_modules是什么npm包管理器和node_modules的关系 2. 如何安装和使用node_modulesnpm安装和使用node_modules的基本命令package.json文件的作用和结构npm包版本号的含义及如何管理包版本 3. 如何发布自己的npm包npm包的结构和规范如何将自己的…
阅读更多...
基于SpringBoot+微信小程序的医院预约叫号小程序

基于SpringBoot+微信小程序的医院预约叫号小程序

✌全网粉丝20W,csdn特邀作者、博客专家、CSDN新星计划导师、java领域优质创作者,博客之星、掘金/华为云/阿里云/InfoQ等平台优质作者、专注于Java技术领域和毕业项目实战✌ &#x1f345;文末获取项目下载方式&#x1f345; 一、项目背景介绍&#xff1a; 该项目是基于uniappWe…
阅读更多...
基于Python的接口自动化测试框架

基于Python的接口自动化测试框架

目录 前言&#xff1a; 项目背景 工具选择 框架思路 第三方库介绍 代码实现 不足之处 前言&#xff1a; Python是一种流行的编程语言&#xff0c;Python的易学性和易用性使得它成为编写接口自动化测试框架的理想选择。在Python中&#xff0c;有许多库可以帮助我们执行HTTP请求…
阅读更多...
淘宝拍照基于端云协同的视频流实时搜索实践

淘宝拍照基于端云协同的视频流实时搜索实践

本文介绍了实时视频流的主体识别场景&#xff0c;未来实时搜将会融合图搜主链路并在XR场景发力&#xff0c;未来的场景我们取名为“元视界”&#xff08;MetaSight&#xff09; 引言 很多熟悉淘宝的用户知道&#xff0c;点击首页搜索框的相机icon&#xff0c;就可以使用淘宝拍照…
阅读更多...
最新文章

Copyright @ 2022~2023