CMU 10-414/714: Deep Learning Systems --hw3

news2024/11/16 1:49:58

实现功能

  1. ndarray.py文件中完成一些python array操作
    • 我们实现的NDArray底层存储就是一个一维向量,只不过会有一些额外的属性(如shape、strides)来表明这个flat array在维度上的分布。底层运算(如加法、矩阵乘法)都是用C++写的,以便达到理论上的最高速度。但像转置、broadcasting、矩阵求子阵等操作都可以在python框架里通过简单的调整array的属性(如strides)来实现
    • 实现的python array功能:这四个功能都不需要重新分配内存,而是在原array的基础上修改shape/strides/etc来实现一些变形(除了第四个__getitem__(),它并不修改array的形状)
      • reshape(new_shape):调用as_strided(new_shape, strides),再调用make(new_shape, strides, device=self.device, handle=self._handle)。(其中as_strided会额外检查一下shape的长度是否等于strides的长度,strides的数值是通过self.compact_strides(new_shape)算出来的)
      • permute(new_axes):调换array的维度顺序。按照new_axes的顺序调换self.strides以及self.shape得到new_shape和new_strides,最后再调用self.as_strided"(new_shape, new_strides)
      • broadcast_to():对于原shape中为1的维度,将其步幅设为0,然后调用as_strided(new_shape, new_strides),然后as_strided调用make(new_shape, new_strides, device=self.device, handle=self._handle)。即重新make,修改其中的shape和stride,但内存位置不变(即handle不变)。至于这里将stride部分设为0,猜测更底层有代码处理这里,但我目前还没找到
      • __getitem__(idxs):这是一个非常好的理解我们所实现的NDArray的例子,前面也说了,底层的array存储就是一个一维向量,是通过shape、stride、offset这些属性来指示该向量的维度的。本例就是在原array的基础上求子矩阵,且不改变其内存位置,就是通过一系列计算得到新的new_offset、new_shape、new_stride,就可以将原本的大矩阵重新解释成一个小矩阵。图示如下
        在这里插入图片描述
        而这里也天然解释了什么是紧凑型(compact)矩阵和非紧凑型矩阵:上述原矩阵就是紧凑型的,切片后形成的蓝底矩阵就是非紧凑型的。
        而这也带出了下一个问题:如何选择原矩阵的一个子矩阵进行修改内容。因为本课程中几乎所有setitem操作都会先调用.compact(),而这会导致从原矩阵中copy子矩阵到一个新的内存空间,这显然不是__setitem__()所期望的,因此要大动干戈。
  2. 【CPU】compact和setitem
    其实这里的实现主要就是根据这个原理:out[cnt++] = in[strides[0]*i + strides[1]*j + strides[2]*k];(这里是按三维矩阵算的,下面的代码通过while循环可以扩展到其他维度)
    • Compact:就是将原来的非紧凑型矩阵a变成紧凑型矩阵out,具体是根据输入的shape、strides、offset(以a为主体确定的),将a中的子矩阵提出来,将各位置的值赋给另一个矩阵out。下面的mode是INDEX_IN,即调用_strided_index_setter(&a, out, shape, strides, offset, INDEX_IN);
      // 从后往前每一个维度,都根据该维度的步长、offset以及循环次数来确定本次要copy矩阵a的哪一个元素
      // 其实就是根据strides、offset来确定一维下的index
      void _strided_index_setter(const AlignedArray* a, AlignedArray* out, std::vector<uint32_t> shape,
                          std::vector<uint32_t> strides, size_t offset, strided_index_mode mode, int val=-1) {
        int depth = shape.size();
        std::vector<uint32_t> loop(depth, 0);
        int cnt = 0;
        while (true) {
          // inner loop
          int index = offset;
          for (int i = 0; i < depth; i++) {
            index += strides[i] * loop[i];
          }
          switch (mode) {
            case INDEX_OUT: out->ptr[index] = a->ptr[cnt++]; break;
            case INDEX_IN: out->ptr[cnt++] = a->ptr[index]; break;
            case SET_VAL: out->ptr[index] = val; break;
          }
      
          // increment
          loop[depth - 1]++;
      
          // carry
          int idx = depth - 1;
          while (loop[idx] == shape[idx]) {
            if (idx == 0) {
              // overflow
              return;
            }
            loop[idx--] = 0;
            loop[idx]++;
          }
        }
      }
      
    • EwiseSetitem(a, out, shape, strides, offset):这里a是紧凑型矩阵,out是一个非紧凑行矩阵,需要将a的各元素的值赋给out的指定位置上(根据shape、strides、offset确定)。只需要复用上面的代码,将mode改为INDEX_OUT,即调用_strided_index_setter(&a, out, shape, strides, offset, INDEX_OUT);
    • ScalarSetitem(size, val, out, shape, strides, offset):out是一个非紧凑型子矩阵,根据shape、strides、offset在out的对应位置将值写为val,即调用_strided_index_setter(nullptr, out, shape, strides, offset, SET_VAL, val);
  3. 【CPU】Elementwise和scalar操作
    这个就非常简单了,根据矩阵的size属性遍历其每个元素即可
  4. 【CPU】Reductions
    这个也很简单,条件是矩阵的底层存储都是一维、行优先的。这类里有两个函数:ReduceSum和ReduceMax,其参数都是原矩阵a、输出矩阵out、reduce_size,其中a.size = out.size * reduce_size。看代码就懂了
    • ReduceMax:
      void ReduceMax(const AlignedArray& a, AlignedArray* out, size_t reduce_size) {
        for (int i = 0; i < out->size; i++) {
          scalar_t max = a.ptr[i * reduce_size];
          for (int j = 0; j < reduce_size; j++) {
            max = std::max(max, a.ptr[i * reduce_size + j]);
          }
          out->ptr[i] = max;
        }
      }
      
    • ReduceSum:
      void ReduceSum(const AlignedArray& a, AlignedArray* out, size_t reduce_size) {
        for (int i = 0; i < out->size; i++) {
          scalar_t sum = 0;
          for (int j = 0; j < reduce_size; j++) {
            sum += a.ptr[i * reduce_size + j];
          }
          out->ptr[i] = sum;
        }
      }
      
  5. 【CPU】矩阵乘
    • 朴素矩阵乘Matmul
      void Matmul(const AlignedArray& a, const AlignedArray& b, AlignedArray* out, uint32_t m, uint32_t n,
                  uint32_t p) {
        for (int i = 0; i < m; i++) {
          for (int j = 0; j < p; j++) {
            out->ptr[i * p + j] = 0;
            for (int k = 0; k < n; k++) {
              out->ptr[i * p + j] += a.ptr[i * n + k] * b.ptr[k * p + j];
            }
          }
        }
      }
      
    • MatmulTiledAlignedDot:前者负责根据分片大小计算目前参与运算的是a、b、out的哪些元素(哪部分block);后者根据前者传进来的block进行计算(a分片与b分片进行矩阵乘法得到out分片
      void MatmulTiled(const AlignedArray& a, const AlignedArray& b, AlignedArray* out, uint32_t m,
                       uint32_t n, uint32_t p) {
        for (int i = 0; i < m * p; i++) out->ptr[i] = 0;
        for (int i = 0; i < m / TILE; i++) {
          for (int j = 0; j < p / TILE; j++) {
            for (int k = 0; k < n / TILE; k++) {
              AlignedDot(&a.ptr[i * n * TILE + k * TILE * TILE], 
                         &b.ptr[k * p * TILE + j * TILE * TILE], 
                         &out->ptr[i * p * TILE + j * TILE * TILE]);
            }
          }
        }
      }
      
      inline void AlignedDot(const float* __restrict__ a,
                             const float* __restrict__ b,
                             float* __restrict__ out) {
        a = (const float*)__builtin_assume_aligned(a, TILE * ELEM_SIZE);
        b = (const float*)__builtin_assume_aligned(b, TILE * ELEM_SIZE);
        out = (float*)__builtin_assume_aligned(out, TILE * ELEM_SIZE);
        
        for (int i = 0; i < TILE; i++) {
          for (int j = 0; j < TILE; j++) {
            for (int k = 0; k < TILE; k++) {
              out[i * TILE + j] += a[i * TILE + k] * b[k * TILE + j];
            }
          }
        }
      }
      
      过程原理图:
      在这里插入图片描述
      这里有一个疑问,a、b虽然每次通过MatmulTiled计算的分片的起始位置是正确的,但是如何保证连续的Tile*Tile分块元素就是如图所示那样的呢。按理说a、b的底层存储应该是连续的行优先,那(按照tile=2算)a[2]、a[3]、b[2]、b[3]应该是按行走下去而不是取下一行啊。待解答…
  6. 【CUDA】compact和setitem
    • compact:
      __device__ size_t index_transform(size_t index, CudaVec shape, CudaVec strides, size_t offset) {
        size_t idxs[MAX_VEC_SIZE];
        size_t cur_size, pre_size = 1;
        // 将给定的线性索引映射回多维数组的索引,即计算index值在各维度上对应是第几个
        // 思路就是从后往前遍历shape,cur_size表示当前处理的维度由几个元素构成
        // index%cur_size/pre_size就表示在当前维度的第几个分量
        // index%cur_size表示是当前维度(只看当前维)的第几个元素,再/pre_size就表示是当前维度的第几块
        for (int i = shape.size - 1; i >= 0; i--) {
          cur_size = pre_size * shape.data[i]; 
          idxs[i] = index % cur_size / pre_size;
          pre_size = cur_size;
        }
        // 根据上述算好的多维数组索引,计算在原非紧凑型矩阵中的线性索引
        size_t comp_idx = offset;
        for (int i = 0; i < shape.size; i++) 
          comp_idx += idxs[i] * strides.data[i];
        return comp_idx;
      }
      
      __global__ void CompactKernel(const scalar_t* a, scalar_t* out, size_t size, CudaVec shape,
                                    CudaVec strides, size_t offset) {
        size_t gid = blockIdx.x * blockDim.x + threadIdx.x;
      
        /// BEGIN YOUR SOLUTION
        if (gid < size)
          out[gid] = a[index_transform(gid, shape, strides, offset)];
        /// END YOUR SOLUTION
      }
      
      void Compact(const CudaArray& a, CudaArray* out, std::vector<uint32_t> shape,
                   std::vector<uint32_t> strides, size_t offset) {
        CudaDims dim = CudaOneDim(out->size);
        CompactKernel<<<dim.grid, dim.block>>>(a.ptr, out->ptr, out->size, VecToCuda(shape),
                                               VecToCuda(strides), offset);
      }
      
      图示如下:
      在这里插入图片描述
    • EWiseSetitem:
      __global__ void EwiseSetitemKernel(const scalar_t* a, scalar_t* out, size_t size, CudaVec shape,
                                         CudaVec strides, size_t offset) {
        size_t gid = blockIdx.x * blockDim.x + threadIdx.x;
        if (gid < size)
          out[index_transform(gid, shape, strides, offset)] = a[gid];
      }
      
      void EwiseSetitem(const CudaArray& a, CudaArray* out, std::vector<uint32_t> shape,
                        std::vector<uint32_t> strides, size_t offset) {
        /// BEGIN YOUR SOLUTION
        CudaDims dim = CudaOneDim(out->size);
        EwiseSetitemKernel<<<dim.grid, dim.block>>>(a.ptr, out->ptr, a.size, VecToCuda(shape),
                                                    VecToCuda(strides), offset);
        /// END YOUR SOLUTION
      }
      
    • ScalarSetitem:
      __global__ void ScalarSetitemKernel(size_t size, scalar_t val, scalar_t* out, CudaVec shape, 
                                          CudaVec strides, size_t offset) {
        size_t gid = blockIdx.x * blockDim.x + threadIdx.x;
        if (gid < size)
          out[index_transform(gid, shape, strides, offset)] = val;
      }
      
      void ScalarSetitem(size_t size, scalar_t val, CudaArray* out, std::vector<uint32_t> shape,
                         std::vector<uint32_t> strides, size_t offset) {
        /// BEGIN YOUR SOLUTION
        CudaDims dim = CudaOneDim(out->size);
        ScalarSetitemKernel<<<dim.grid, dim.block>>>(size, val, out->ptr, VecToCuda(shape),
                                                     VecToCuda(strides), offset);
        /// END YOUR SOLUTION
      }
      
      • 其实总结一下,CPU版的将紧凑小矩阵的index转换成非紧凑大矩阵的index,是在一个loop中实现的,并在loop中找到index后就完成了copy工作;对于GPU版来说,是将copy工作分配给各个线程,因此若要让每个线程都能正确copy,还需要每个线程根据自己分配到的紧凑小矩阵的index,计算得到非紧凑大矩阵的index。即每个线程完成CPU版中一个loop的操作(但不需要后续的检查)。
  7. 【CUDA】Elementwise和scalar操作
    和CPU版本的一样,也很简单。且这part非常能体现CUDA的并行、高效特性。举个例子:
    __global__ void EwiseMulKernel(const scalar_t* a, const scalar_t* b, scalar_t* out, size_t size){
    	size_t gid = blockInx.x * blockDim.x + threadIdx.x;
    	if (gid<size){
    		out[gid] = a[gid] * b[gid];
    	}
    }
    void EwiseMul(const CudaArray &a, const CudaArray &b, CudaArray* out){
    	CudaDims dim = CudaOneDim(out->size);
    	EwiseMulKernel<<<dim.grid, dim.block>>>(a.ptr, b.ptr, out->ptr, out->size);
    }
    
  8. 【CUDA】Reductions
    • ReduceMax
      __global__ void ReduceMaxKernel(const scalar_t* a, scalar_t* out, size_t reduce_size, size_t size){
        size_t gid = blockIdx.x * blockDim.x + threadIdx.x;
        if (gid<size){
          size_t offset = gid * reduce_size;
          scalar_t reduce_max = a[offset];
          for (int i=1; i<reduce_size; i++){
            reduce_max = max(reduce_max, a[offset+i]);
          }
          out[gid] = reduce_max;
        }
      }
      
      void ReduceMax(const CudaArray& a, CudaArray* out, size_t reduce_size) {
        CudaDims dim = CudaOneDim(out->size);
        ReduceMaxKernel<<<dim.grid, dim.block>>>(a.ptr, out->ptr, reduce_size, out->size);
      }
      
    • ReduceSum
  9. 【CUDA】矩阵乘
    使用朴素矩阵乘即可,一个线程负责out矩阵的一个元素
    __global__ void MatmulKernel(const scalar_t* a, const scalar_t* b, scalar_t* out, uint32_t M,
                                 uint32_t N, uint32_t P) {
      size_t i = blockIdx.x * blockDim.x + threadIdx.x; 
      size_t j = blockIdx.y * blockDim.y + threadIdx.y;
      if (i < M && j < P) {
        out[i * P + j] = 0;
        for (int k = 0; k < N; k++) {
          out[i * P + j] += a[i * N + k] * b[k * P + j];
        }
      }
    }
    void Matmul(const CudaArray& a, const CudaArray& b, CudaArray* out, uint32_t M, uint32_t N,
                uint32_t P) {
      dim3 grid(BASE_THREAD_NUM, BASE_THREAD_NUM, 1);
      dim3 block((M + BASE_THREAD_NUM - 1) / BASE_THREAD_NUM, (P + BASE_THREAD_NUM - 1) / BASE_THREAD_NUM, 1);
      MatmulKernel<<<grid, block>>>(a.ptr, b.ptr, out->ptr, M, N, P);
    }
    
    线程分布如下:
    在这里插入图片描述

补充知识

一、CUDA相关代码

在这里插入图片描述
如上图的结构,首先定义Cuda的维度:

struct CudaDims {
  dim3 block, grid;
};

根据数据数量来判定需要多少个block、多少个thread(都是一个维度下的)

CudaDims CudaOneDim(size_t size) {
  /**
   * Utility function to get cuda dimensions for 1D call
   */
  CudaDims dim;
  size_t num_blocks = (size + BASE_THREAD_NUM - 1) / BASE_THREAD_NUM;
  dim.block = dim3(BASE_THREAD_NUM, 1, 1);  // 一个block里的线程
  dim.grid = dim3(num_blocks, 1, 1);  // 一个grid里的block
  return dim;
}

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

如若内容造成侵权/违法违规/事实不符,请联系多彩编程网进行投诉反馈,一经查实,立即删除!

相关文章

网络安全的几个关键领域

网络安全是一个复杂且多维度的领域&#xff0c;涵盖了多个关键领域&#xff0c;涉及到信息保护、网络防护、应用安全、用户教育以及物理安全等多个方面。这些关键领域相互交织&#xff0c;共同构成了网络安全这一宏大且细致入微的领域。 今天德迅云安全就分享下网络安全的几个…

android studio的布局没有提示之SDK不匹配

我新建了一个项目&#xff0c;然后突然发现布局没有提示了&#xff1a; 我看了下我的build.gradle 我直接修改compileSdkVersion为30就能正常使用了

【贪心算法】Leetcode 55. 跳跃游戏 45. 跳跃游戏 I

【贪心算法】Leetcode 55. 跳跃游戏 45. 跳跃游戏 II Leetcode 55. 跳跃游戏解法1 贪心 Leetcode 45. 跳跃游戏I解法 贪心 Leetcode 55. 跳跃游戏 ---------------&#x1f388;&#x1f388;55. 跳跃游戏 题目链接&#x1f388;&#x1f388;------------------- 解法1 贪心…

【计算机网络篇】物理层(2)传输方式

文章目录 &#x1f354;传输方式⭐串行传输和并行传输⭐同步传输和异步传输&#x1f388;同步传输&#x1f388;异步传输 ⭐单向通信&#xff0c;双向交替通信和双向同时通信 &#x1f354;传输方式 在物理层中&#xff0c;传输方式是指数据在传输介质中传输的方式和方法。常见…

《论文阅读》EmpDG:多分辨率交互式移情对话生成 COLING 2020

《论文阅读》EmpDG:多分辨率交互式移情对话生成 COLING 2020 前言简介模型架构共情生成器交互鉴别器损失函数前言 亲身阅读感受分享,细节画图解释,再也不用担心看不懂论文啦~ 无抄袭,无复制,纯手工敲击键盘~ 今天为大家带来的是《EmpDG: Multi-resolution Interactive E…

牛客题霸-SQL进阶篇(刷题记录一)

本文基于前段时间学习总结的 MySQL 相关的查询语法&#xff0c;在牛客网找了相应的 MySQL 题目进行练习&#xff0c;以便加强对于 MySQL 查询语法的理解和应用。 由于涉及到的数据库表较多&#xff0c;因此本文不再展示&#xff0c;只提供 MySQL 代码与示例输出。 部分题目因…

递归算法c++

主页:(*∇&#xff40;*) 咦,又好了~ xiaocr_blog 算法概述&#xff1a;递归算法是一种直接或者间接调用自身函数或者方法的算法。说简单了就是程序自身的调用。 算法实质&#xff1a;递归算法就是将原问题不断分解为规模缩小的子问题&#xff0c;然后递归调用方法来表示问题的…

esp8266WIFI模块教程:ATK-ESP8266——TCP网络通讯和服务器连接

前言 ATK-MW8266D 模块是正点原子推出的一款高性能 UART-WIFI&#xff08;串口-无线&#xff09;模块&#xff0c;ATK-MW8266D 模块板载了正点原子公司自主研发的 ATK-ESP-01 模块&#xff0c;该模块通过 FCC、CE 认证&#xff0c;可直接用于出口欧美地区的产品。   ATK-MW82…

springboot酒店管理系统 论文【源码】

springboot酒店管理系统开发说明 开发语言&#xff1a;Java 框架&#xff1a;springboot JDK版本&#xff1a;JDK1.8 服务器&#xff1a;tomcat7 数据库&#xff1a;mysql 5.7&#xff08;一定要5.7版本&#xff09; 数据库工具&#xff1a;Navicat11 开发软件&#xff1…

【SQL】1934. 确认率(IFNULL(); IF())

题目描述 leetcode题目&#xff1a;1934. 确认率 Code select Signups.user_id,ifnull(confirmation_rate1, 0) as confirmation_rate from Signups left join(select user_id,round(sum(if(actionconfirmed, 1, 0))/count(*), 2) as confirmation_rate1from Confirmationsgr…

蓝桥杯-24点-搜索

题目 思路 --暴力递归全组合的方法。只有4个数&#xff0c;4种计算方式&#xff0c;共有4 * 3 * 2 * 1 * 4种不同的情况&#xff0c;可以写递归来实现。 --每次计算都是两个数之间的运算&#xff0c;因此4个数需要3次计算&#xff0c;第一次计算前有4个数&#xff0c;第二次有…

Python:filter过滤器

filter() 是 Python 中的一个内置函数&#xff0c;用于过滤序列&#xff0c;过滤掉不符合条件的元素&#xff0c;返回由符合条件元素组成的新列表。该函数接收两个参数&#xff0c;一个是函数&#xff0c;一个是序列&#xff0c;序列的每个元素作为参数传递给函数进行判定&…

耐腐蚀高纯特氟龙塑料量瓶进口聚四氟乙烯材质PFA容量瓶

PFA容量瓶&#xff0c;也叫特氟龙量瓶&#xff0c;是用于配制标准浓度溶液的精确实验室器皿&#xff0c;是有着细长颈、梨形肚的耐强腐蚀平底塑料瓶&#xff0c;颈上有标线&#xff0c;常用来直接配制标准溶液和准确稀释溶液以及制备样品溶液。因其有着不易碎、材质纯净、化学稳…

Android studio开发中Virtual Device模拟器的设置和屏幕错位等问题

Android SDK开发中Virtual Device模拟器的设置和使用 本文介绍android studio2023 3.1.13版本中模拟器的设置和在cordova开发中的运行方法 对于老版android studioAVD模拟器的使用&#xff0c;参见&#xff1a;Android SDK手机应用开发中第三方模拟器、真机运行方法以及AVD模拟…

camelot pdf提取表格实践(记录)

参考&#xff1a; 巧用Python的camelot库批量提取PDF发票信息 关于文本pdf的表格抽取 AttributeError: module ‘camelot‘ has no attribute ‘read_pdf‘及类似问题解决办法 camelot 参数 https://blog.csdn.net/INTSIG/article/details/123000010 报错解决&#xff1a; Mod…

简析|创业老隋分享的人力RPO项目如何?

在当今社会&#xff0c;创业热潮席卷而来&#xff0c;各类项目层出不穷。近日&#xff0c;创业老隋分享的人力RPO项目引起了广泛关注。那么&#xff0c;这个项目究竟如何呢?是否靠谱?经过深入了解和分析&#xff0c;我认为这个项目是靠谱的。 首先&#xff0c;从项目的背景和…

2.1(TCP)

TCP—传输控制协议 是一种面向连接的可靠传输协议。可靠、有序、无丢弃和不重复。 特点&#xff1a; TCP是面向连接&#xff08;虚连接&#xff09;的传输层协议每一条TCP连接有且只能有两个端点。可靠、有序、无丢弃和不重复。TCP协议提供全双工通讯。 发送缓存 存放发送方…

Linux--Docker容器引擎

1、Docker是什么 Docker是在Linux容器里运行应用的开源工具&#xff0c;是一种轻量级的 “虚拟机”。诞生于2013年&#xff0c;最 初的发起者是dotCloud公司。后来改名为Docker Inc&#xff0c;专注于Docker相关技术和产品的开发。Docker 项目目前已经加人啊Linux基金会&#x…

迅雷网盘推广拉新怎么申请授权怎么赚钱?

迅雷网盘作为国内知名的云存储服务提供商&#xff0c;拥有庞大的用户群体和稳定的服务质量。为了进一步扩大市场份额&#xff0c;迅雷网盘推出了拉新项目&#xff0c;通过给予推广者一定的佣金&#xff0c;鼓励他们向更多的人推广迅雷网盘。 网盘功能强大&#xff1a;作为国内比…

【送外卖的你】送外卖对程序员的职业发展有影响吗

送外卖对程序员的职业发展可能会有一些影响&#xff0c;具体取决于个人情况和职业目标。以下是一些可能的影响&#xff1a; 1. 时间管理&#xff1a;送外卖可能需要花费较多的时间和精力&#xff0c;这可能会对程序员的工作时间和项目进度产生一定的影响。需要合理安排时间&…