[CUDA 学习笔记] 矩阵转置算子优化

news2025/9/22 1:40:30

矩阵转置算子优化

矩阵转置是一种基础的矩阵操作, 即将二维矩阵的行列进行反转.
本文主要围绕行主序的二维单精度矩阵的转置考虑相关的优化.

以下 kernel 笔者均是在 NVIDIA V100 (7.0 算力) 上进行测试的, 且选择矩阵的行列维度大小为 M=2300 N=1500.

Version 0. 朴素实现

__global__ void mat_transpose_kernel_v0(const float* idata, float* odata, int M, int N) {
    const int x = blockIdx.x * blockDim.x + threadIdx.x;
    const int y = blockIdx.y * blockDim.y + threadIdx.y;
    
    if (y < M && x < N) {
        odata[x * M + y] = idata[y * N + x];
    }
}


void mat_transpose_v0(const float* idata, float* odata, int M, int N) {
    constexpr int BLOCK_SZ = 16;
    dim3 block(BLOCK_SZ, BLOCK_SZ);
    dim3 grid((N + BLOCK_SZ - 1) / BLOCK_SZ, (M + BLOCK_SZ - 1) / BLOCK_SZ));
    mat_transpose_kernel_v0<<<grid, block>>>(idata, odata, M, N);
}

矩阵转置的朴素实现非常直观, 思路即使用二维的线程/线程块排布, 让每个线程负责一个矩阵元素的转置. 实现上, 只需要将矩阵的行列索引 x y 进行反转即可.
需要注意的是 grid 和 block 的中维度设置与多维数组中的表示是相反的, 即 grid.x 应该对应 N 维度, grid.y 应该对应 M 维度.

Version用时(us)内存带宽(GB/s)带宽利用率(%)加速比
v084.90399.9256.38

结合矩阵转置的逻辑以及 Nsight Compute 容易判断出, 矩阵转置本身是一个 memory-bound 的 kernel, 因为其核心是完成矩阵内存排布的转换, 这个过程基本不涉及计算, 因此对该 kernel 优化很重要的一点就是提高访存性能.
朴素实现直接操作矩阵所在的 GMEM, 直观看来, 矩阵转置不会涉及数据的重用, 直接操作 GMEM 本身没有问题, 但此时应该注意 GMEM 的访存特性, 其中很重要的即 GMEM 的访存合并, 即连续线程访问的 GMEM 中的数据地址是连续的, 可以将多个线程的内存访问合并为一个(或多个)内存访问, 从而减少访存次数, 提高带宽利用率.
在 Version 0 的 kernel 中, 容易看出读取时 idata[y * N + x] 是访存合并的, 因为连续线程对应的 x 是连续的, 即访问矩阵同一行连续的列; 但是写入时 odata[x * M + y] 并不是访存合并的, 因为转置后连续线程写入的是同一列连续的行, 但由于内存布局是行主序的, 因此此时每个线程访问的地址实际上并不连续, 地址差 N, 因此对 GMEM 访存性能有很大影响.

Version 1. 利用共享内存合并访存

template <int BLOCK_SZ>
__global__ void mat_transpose_kernel_v1(const float* idata, float* odata, int M, int N) {
    const int bx = blockIdx.x, by = blockIdx.y;
    const int tx = threadIdx.x, ty = threadIdx.y;

    __shared__ float sdata[BLOCK_SZ][BLOCK_SZ];
    
    int x = bx * BLOCK_SZ + tx;
    int y = by * BLOCK_SZ + ty;

    if (y < M && x < N) {
        sdata[ty][tx] = idata[y * N + x];
    }
    __syncthreads();

    x = by * BLOCK_SZ + tx;
    y = bx * BLOCK_SZ + ty;
    if (y < N && x < M) {
        odata[y * M + x] = sdata[tx][ty];
    }
}

void mat_transpose_v1(const float* idata, float* odata, int M, int N) {
    constexpr int BLOCK_SZ = 16;
    dim3 block(BLOCK_SZ, BLOCK_SZ);
    dim3 grid(Ceil(N, BLOCK_SZ), Ceil(M, BLOCK_SZ));
    mat_transpose_kernel_v1<BLOCK_SZ><<<grid, block>>>(idata, odata, M, N);
}

Version 0 的 kernel 存在的问题是写入 GMEM 时访存不合并, 因此需要一种方式让写入 GMEM 时仍然保持线程的访存连续, 在 Version 1 中, 便使用了 SMEM 用来中转来实现访存合并.
在这里插入图片描述
Version 1 的核心思想可以使用上图进行表示, 中间的 “tile” 即可理解为存在 SMEM 的数据分片.
在读取矩阵阶段, 操作与 Version 0 一致, 区别在于将数据直接写入 SMEM 中, 对应上图橙色部分. 接着通过设置 x = by * BLOCK_SZ + tx; y = bx * BLOCK_SZ + ty; 两条语句进行了索引的重计算, 进行了线程块索引 bx 和 by 交换, 对应上图右上角的数据分片转置后成为了左下角的数据分片. 由于此时 tx 和 ty 并没有交换, 因此按照 odata[y * M + x] 写入 GMEM 时, 访存是合并的, 但需要读取 SMEM 时 tx 与 ty 进行交换, 实现数据分片内的转置, 对应上图绿色部分.

Version用时(us)内存带宽(GB/s)带宽利用率(%)加速比
v084.90399.9256.38
v147.49610.5771.071.79

可以看到, Version 1 相比于 Version 0 性能有了很大提升. 如下图所示, 通过 Nsight Compute 也能看到 Version 1 (右) 比 Version 0 (左) 在读取写入 GMEM 的带宽上都有所提升.
在这里插入图片描述 在这里插入图片描述
在 Version 1 中, 引入了对 SMEM 的访存, 因此需要特别关注 bank conflict 的问题. 对于 SMEM 的写入, sdata[ty][tx], 由于 BLOCK_SZ 为 16, 32 个线程负责 SMEM 矩阵分片的两行 32 个元素, 对应 32 个 bank, 因此没有 bank conflict. 而对于 SMEM 的读取, sdata[tx][ty], 由于是按列读取 SMEM 的, threadIdx 差 1 的线程访问的数据差 BLOCK_SZ, 即导致threadIdx 差 2 的线程访问的数据落到同一 bank 的不同地址, 从而造成 16 路的 bank conflict.
因此, Version 1 中读取 SMEM 引起的 bank conflict 影响 SMEM 的访存效率, 进而影响 kernel 的性能.

Version 2. 利用 padding 解决 bank conflict

template <int BLOCK_SZ>
__global__ void mat_transpose_kernel_v2(const float* idata, float* odata, int M, int N) {
    const int bx = blockIdx.x, by = blockIdx.y;
    const int tx = threadIdx.x, ty = threadIdx.y;

    __shared__ float sdata[BLOCK_SZ][BLOCK_SZ+1];    // padding
    
    int x = bx * BLOCK_SZ + tx;
    int y = by * BLOCK_SZ + ty;

    if (y < M && x < N) {
        sdata[ty][tx] = idata[y * N + x];
    }
    __syncthreads();

    x = by * BLOCK_SZ + tx;
    y = bx * BLOCK_SZ + ty;
    if (y < N && x < M) {
        odata[y * M + x] = sdata[tx][ty];
    }
}

void mat_transpose_v2(const float* idata, float* odata, int M, int N) {
    constexpr int BLOCK_SZ = 16;
    dim3 block(BLOCK_SZ, BLOCK_SZ);
    dim3 grid(Ceil(N, BLOCK_SZ), Ceil(M, BLOCK_SZ));
    mat_transpose_kernel_v2<BLOCK_SZ><<<grid, block>>>(idata, odata, M, N);
}

Version 2 的代码相比于 Version 1 仅在 SMEM 内存分配时进行了变动, 将大小改为了 sdata[BLOCK_SZ][BLOCK_SZ+1], 即列维度上加入了 1 元素大小的 padding.
此时, 对于读取 SMEM 的 sdata[tx][ty], threadIdx 差 1 的线程访问的数据差 BLOCK_SZ+1, 即 17, 由于 17 与 32 互质, 因此不会有 bank conflict. 值得一提的是, 对于写入 SMEM 的 sdata[ty][tx], 由于有 1 个 padding, warp 中 lane 31 与 lane 0 访问的元素恰好差 31+1=32 个元素, 会有 1 个 bank conflict.
整体上, Version 2 通过 padding 的方法有效避免了读取 SMEM 时的 bank conflict.

Version用时(us)内存带宽(GB/s)带宽利用率(%)加速比
v084.90399.9256.38
v147.49610.5771.071.79
v244.38627.3172.471.91

Version 2 相比 Version 1性能有了一定的提升. 如下图所示, 通过 Nsight Compute 也能看到 Version 2 (右) 比 Version 1 (左) 的 bank conflict 总数明显下降, 其中读取 SMEM 时的 bank conflict (第一行) 大幅降低, 而写入 SMEM 时的 bank conflict (第二行) 有一定上升, 也与上文分析的相匹配.
在这里插入图片描述 在这里插入图片描述

Version 3. 增加每个线程处理的元素个数

template <int BLOCK_SZ, int NUM_PER_THREAD>
__global__ void mat_transpose_kernel_v3(const float* idata, float* odata, int M, int N) {
    const int bx = blockIdx.x, by = blockIdx.y;
    const int tx = threadIdx.x, ty = threadIdx.y;

    __shared__ float sdata[BLOCK_SZ][BLOCK_SZ+1];
    
    int x = bx * BLOCK_SZ + tx;
    int y = by * BLOCK_SZ + ty;

    constexpr int ROW_STRIDE = BLOCK_SZ / NUM_PER_THREAD;

    if (x < N) {
        #pragma unroll
        for (int y_off = 0; y_off < BLOCK_SZ; y_off += ROW_STRIDE) {
            if (y + y_off < M) {
                sdata[ty + y_off][tx] = idata[(y + y_off) * N + x]; 
            }
        }
    }
    __syncthreads();

    x = by * BLOCK_SZ + tx;
    y = bx * BLOCK_SZ + ty;
    if (x < M) {
        for (int y_off = 0; y_off < BLOCK_SZ; y_off += ROW_STRIDE) {
            if (y + y_off < N) {
                odata[(y + y_off) * M + x] = sdata[tx][ty + y_off];
            }
        }
    }
}

void mat_transpose_v3(const float* idata, float* odata, int M, int N) {
    constexpr int BLOCK_SZ = 32;
    constexpr int NUM_PER_THREAD = 4;
    dim3 block(BLOCK_SZ, BLOCK_SZ/NUM_PER_THREAD);
    dim3 grid(Ceil(N, BLOCK_SZ), Ceil(M, BLOCK_SZ));
    mat_transpose_kernel_v3<BLOCK_SZ, NUM_PER_THREAD><<<grid, block>>>(idata, odata, M, N);
}

Version 3 相比于 Version 2, 增加了每个线程处理的元素个数, 即由先前的每个线程处理 1 个元素的转置, 变为处理 NUM_PER_THREAD 个元素的转置. 该实现主要是参考了 英伟达的技术博客.
在实现上, 同样保持原本 256 线程的线程块大小, 设置每个线程处理 4 个元素, 则每个线程块数据分片的大小调整为 32×32, 而线程块的线程采取 8×32 的二维排布, 因此需要在行维度上需要迭代 4 次完成转置.

考虑 Version 3 相比于 Version 2 的优势, 主要是在保持线程块中线程数量不变的情况下, 处理的线程块数据分片大小变大, 这样会减少线程网格中启动的线程块数量, 而增大了每个线程的计算强度; 此外, 由于 BLOCK_SZ 变为 32, Version 2 中写入 SMEM 的 1 个 bank conflict 也可以被避免.
这让笔者想到了 Reduce 算子中也会考虑增加每个线程处理的元素来提高性能, 笔者主观的感觉是对于这种计算强度比较低的 kernel, 增加线程处理的元素个数即计算强度, 一定程度上能增大 GPU 中计算与访存的掩盖, 并配合循环展开提高指令级并行; 此外, 由于线程块数量的减少, 能在相对少的 wave 中完成计算, 减少 GPU 的线程块调度上可能也会带来性能的收益.

Version用时(us)内存带宽(GB/s)带宽利用率(%)加速比
v084.90399.9256.38
v147.49610.5771.071.79
v244.38627.3172.471.91
v339.68699.0581.432.14

通过测试可以看出, Version 3 相比 Version 2 又有了一定的性能提升, 内存带宽也有了进一步的提高.

Version 3.5 向量化读取

#define FETCH_CFLOAT4(p) (reinterpret_cast<const float4*>(&(p))[0])
#define FETCH_FLOAT4(p) (reinterpret_cast<float4*>(&(p))[0])

template <int BLOCK_SZ>
__global__ void mat_transpose_kernel_v3_5(const float* idata, float* odata, int M, int N) {
    const int bx = blockIdx.x, by = blockIdx.y;
    const int tx = threadIdx.x, ty = threadIdx.y;

    __shared__ float sdata[BLOCK_SZ][BLOCK_SZ];
    
    int x = bx * BLOCK_SZ + tx * 4;
    int y = by * BLOCK_SZ + ty;

    if (x < N && y < M) {
        FETCH_FLOAT4(sdata[ty][tx * 4]) = FETCH_CFLOAT4(idata[y * N + x]);
    }
    __syncthreads();

    x = by * BLOCK_SZ + tx * 4;
    y = bx * BLOCK_SZ + ty;
    float tmp[4];
    if (x < M && y < N) {
        #pragma unroll
        for (int i = 0; i < 4; ++i) {
            tmp[i] = sdata[tx * 4 + i][ty];
        }
        FETCH_FLOAT4(odata[y * M + x]) = FETCH_FLOAT4(tmp);
    }
}

void mat_transpose_v3_5(const float* idata, float* odata, int M, int N) {
    constexpr int BLOCK_SZ = 32;
    dim3 block(BLOCK_SZ / 4, BLOCK_SZ);
    dim3 grid(Ceil(N, BLOCK_SZ), Ceil(M, BLOCK_SZ));
    mat_transpose_kernel_v3_5<BLOCK_SZ><<<grid, block>>>(idata, odata, M, N);
}

Version 3 是通过在行维度上迭代多次实现了每个线程处理多个元素的目的. 这里笔者尝试使用向量化访存的形式在列维度上让每个线程负责连续的多个元素.

在实现上, BLOCK_SZ 仍然是 32, 即线程块数据分片的大小保持 32×32, 而线程块的线程采取 32×8 的二维排布, 这是因为列维度上使用向量化访存一次读取 4 个元素.
但是, 向量化访存也带了一些问题, 首先向量化访存需要地址对齐, 因此不能使用 Version 2 中的 padding 方式, 因为 padding 的 1 个元素会导致访问元素并不是 4 个元素地址对齐的; 另一方面, 在从 SMEM 转置写入 GMEM 输出时, 是按列读取 SMEM 的, 无法直接使用向量化访存, 因此需要逐元素读取到寄存器后再向量化访存写入 GMEM.

Version用时(us)内存带宽(GB/s)带宽利用率(%)加速比
v084.90399.9256.38
v147.49610.5771.071.79
v244.38627.3172.471.91
v339.68699.0581.432.14
v3.542.46679.8082.642.00

通过测试可以看出, Version 3.5 的性能相比 Version 2 有一定提升, 但并不如 Version 3., 受益于向量化访存, 带宽相比于 Version 2 也有一定提升, 但由于无法 padding 避免 bank conflict, 也导致其性能并不如 Version 3.

Version 4. 减少条件分支

template <int BLOCK_SZ, int NUM_PER_THREAD>
__global__ void mat_transpose_kernel_v4(const float* idata, float* odata, int M, int N) {
    const int bx = blockIdx.x, by = blockIdx.y;
    const int tx = threadIdx.x, ty = threadIdx.y;

    __shared__ float sdata[BLOCK_SZ][BLOCK_SZ+1];
    
    int x = bx * BLOCK_SZ + tx;
    int y = by * BLOCK_SZ + ty;

    constexpr int ROW_STRIDE = BLOCK_SZ / NUM_PER_THREAD;

    if (x < N) {
        if (y + BLOCK_SZ <= M) {
            #pragma unroll
            for (int y_off = 0; y_off < BLOCK_SZ; y_off += ROW_STRIDE) {
                sdata[ty + y_off][tx] = idata[(y + y_off) * N + x]; 
            }
        } else {
            for (int y_off = 0; y_off < BLOCK_SZ; y_off += ROW_STRIDE) {
                if (ty + y_off < M) {
                    sdata[ty + y_off][tx] = idata[(y + y_off) * N + x];
                }
            }
        }

    }
    __syncthreads();

    x = by * BLOCK_SZ + tx;
    y = bx * BLOCK_SZ + ty;
    if (x < M) {
        if (y + BLOCK_SZ <= N) {
            #pragma unroll
            for (int y_off = 0; y_off < BLOCK_SZ; y_off += ROW_STRIDE) {
                odata[(y + y_off) * M + x] = sdata[tx][ty + y_off];
            }
        } else {
            for (int y_off = 0; y_off < BLOCK_SZ; y_off += ROW_STRIDE) {
                if (y + y_off < N) {
                    odata[(y + y_off) * M + x] = sdata[tx][ty + y_off];
                }
            }
        }
    }
}

void mat_transpose_v4(const float* idata, float* odata, int M, int N) {
    constexpr int BLOCK_SZ = 32;
    constexpr int NUM_PER_THREAD = 4;
    dim3 block(BLOCK_SZ, BLOCK_SZ/NUM_PER_THREAD);
    dim3 grid(Ceil(N, BLOCK_SZ), Ceil(M, BLOCK_SZ));
    mat_transpose_kernel_v4<BLOCK_SZ, NUM_PER_THREAD><<<grid, block>>>(idata, odata, M, N);
}

Version 4 参考了 CUTLASS 中 N C H W NCHW NCHW 与 N H W C NHWC NHWC 格式转换的 kernel, 该 kernel 本质上就是 C C C 与 H W HW HW 两个维度的矩阵转置.
在实现上, Version 4 的 kernel 与 Version 3 的核心思路是一致的, 仍然是每个线程需要行维度上迭代多次. 不同的是, 该实现对分支语句的判断进行了更加细致的处理. 具体而言, 在 Version 3 中, 每个线程的每次行迭代都要进行 y + y_off 的判断, 即矩阵元素的范围检查, 但考虑大部分线程块的数据分片都是在矩阵内部的, 即不会越界, 那么执行多次范围检查就会比较多余. 在 Version 4 中, 会先检查 y + BLOCK_SZ <= M, 即如果该线程块的数据分片并不在矩阵边缘, 那么每次迭代时便无需再进行越界检查; 反之, 若数据分片位于矩阵边缘, 则需要和 Version 3 一样每次迭代时越界检查.
整体看来, 该实现虽然增加了代码复杂程度, 但会简化大部分线程块中迭代时的逻辑, 减少冗余的条件分支的判断次数.

Version用时(us)内存带宽(GB/s)带宽利用率(%)加速比
v084.90399.9256.38
v147.49610.5771.071.79
v244.38627.3172.471.91
v339.68699.0581.432.14
v3.542.46679.8082.642.00
v440.42686.5080.252.10

有意思的是, 在实际测试时, Version 4 的性能与 Version 3 很接近, 但与 Version 3 仍有略微的性能差距, 且带宽上也并没有 Version 3 高.

参考文献

  • An Efficient Matrix Transpose in CUDA C/C++ | NVIDIA Technical Blog
  • CUDA 编程入门之矩阵转置 - 知乎
  • CUDA学习(二)矩阵转置及优化(合并访问、共享内存、bank conflict) - 知乎
  • NVIDIA/cutlass - device_nchw_to_nhwc.h

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

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

相关文章

流量分组新增两大新规则；Network SDK更新618大促版本；综合报表支持实时新用户指标 | TopOn产品更新

流量分组新增两大新规则；Network SDK更新618大促版本；综合报表支持实时新用户指标 | TopOn产品更新

「TopPro 每月产品速递」是由TopOn最新推出的产品专栏&#xff0c;将会以月为周期梳理TopOn最新产品动态&#xff0c;致力于为互联网从业者提供优质服务&#xff0c;引领行业产品发展。 TopPro | 四月产品速递 2023.04.01-04.27 01 流量分组新增两大新规则 // 功能描述 *…
阅读更多...
面向AI编程，AI可以为我们做哪些事情

面向AI编程，AI可以为我们做哪些事情

本来这篇文章是2023-10月发出的&#xff0c;放在草稿箱比较久了。今天重新捡起来发下。内容很长&#xff0c;很干。希望对大家有启发&#xff0c;编程路上提升效率。 背景 基本上以前我们出了bug都是百度&#xff0c;但随着AI的出现&#xff0c;对标百度给出的答案。发现AI实在…
阅读更多...
python爬虫-------JsonPath（第十九天）

python爬虫-------JsonPath（第十九天）

&#x1f388;&#x1f388;作者主页&#xff1a; 喔的嘛呀&#x1f388;&#x1f388; &#x1f388;&#x1f388;所属专栏&#xff1a;python爬虫学习&#x1f388;&#x1f388; ✨✨谢谢大家捧场&#xff0c;祝屏幕前的小伙伴们每天都有好运相伴左右&#xff0c;一定要天天…
阅读更多...
Pytorch中nn.Linear使用方法

Pytorch中nn.Linear使用方法

nn.Linear定义一个神经网络的线性层&#xff1a; torch.nn.Linear(in_features, # 输入的神经元个数out_features, # 输出神经元个数biasTrue # 是否包含偏置)nn.Linear其实就是对输入&#xff08;n表示样本数量&#xff0c;i表示样本特…
阅读更多...
国产低代码工具，轻松搞定数据迁移

国产低代码工具，轻松搞定数据迁移

在日常的业务系统升级或者数据维护过程中&#xff0c;数据迁移是各个企业用户不得不面临的问题&#xff0c;尤其是数据迁移过程中要保障数据完整性、统一性和及时性&#xff0c;同时也需要注意源数据中的数据质量问题&#xff0c;比如缺失、无效、错误等问题&#xff0c;需要在…
阅读更多...
Kubernetes中安装部署Nacos集群

Kubernetes中安装部署Nacos集群

目录 1、Nacos安装包的准备 1.1 下载安装包 1.2 解压安装包 1.3 修改配置文件 application.properties 1.4 bin目录下创建 docker-startup.sh 1.5 将nacos-server-1.2.1目录打包成nacos-server-1.2.1.tar.gz 2、 nacos镜像制作 2.1 Dockerfile文件编写 2.2 制作镜像…
阅读更多...
单片机入门还能从51开始吗？

单片机入门还能从51开始吗？

选择从51单片机开始入门还是直接学习基于ARM核或RISC核的单片机&#xff0c;取决于学习目标、项目需求以及个人兴趣。每种单片机都有其特定的优势和应用场景&#xff0c;了解它们的特点可以帮助你做出更合适的选择。 首先&#xff0c;我们说一下51单片机的优势&#xff1a; 成熟…
阅读更多...
外包干了17天，技术倒退明显

外包干了17天，技术倒退明显

先说情况&#xff0c;大专毕业&#xff0c;18年通过校招进入湖南某软件公司&#xff0c;干了接近6年的功能测试&#xff0c;今年年初&#xff0c;感觉自己不能够在这样下去了&#xff0c;长时间呆在一个舒适的环境会让一个人堕落&#xff01; 而我已经在一个企业干了四年的功能…
阅读更多...
【石上星光】context，go的上下文存储并发控制之道

【石上星光】context，go的上下文存储并发控制之道

目录 1 引言2 What&#xff1f;3 How&#xff1f; 3.1 用法一、上下文数据存储3.2 用法二、并发控制 3.2.1 场景1 主动取消3.2.2 场景2 超时取消 3.3 用法三、创建一个空Context&#xff08;emptyCtx&#xff09; 4 Why&#xff1f; 4.1 go中的上下文思想 4.1.1 上下文是什么…
阅读更多...
技术小课堂：100％CC防护是怎么实现的？

技术小课堂：100％CC防护是怎么实现的？

大家好&#xff0c;今天我们深入探讨的是如何有效地实现CC攻击的100%防护&#xff0c;以及传统防护手段存在的局限性和我们的定制化解决方案的优势。 传统的CC防护措施通常依赖于全局性的访问频率控制或在防火墙级别设置固定的访问次数限制。这种方式看似简单直接&#xff0c;…
阅读更多...
安全大脑与盲人摸象

安全大脑与盲人摸象

21世纪是数字科技和数字经济爆发的时代&#xff0c;互联网正从网状结构向类脑模型进行进化&#xff0c;出现了结构和覆盖范围庞大&#xff0c;能够适应不同技术环境、经济场景&#xff0c;跨地域、跨行业的类脑复杂巨型系统。如腾讯、Facebook等社交网络具备的神经网络特征&…
阅读更多...
[方案实操|数据技术]数据要素十大创新模式(1)：基于区块链的多模态数据交易服务平台

[方案实操|数据技术]数据要素十大创新模式(1)：基于区块链的多模态数据交易服务平台

“ 区块链以其公开共享、去中心化、不可篡改、可追溯和不可抵赖等优势&#xff0c;吸引了包括金融业、医疗业和政府部门等众多利益相关方的极大兴趣&#xff0c;被认为是解决数据安全交换问题的合适方案。” 武汉东湖大数据科技股份有限公司凭借基于区块链的多模态数据交易服务…
阅读更多...
交换机的基本原理与配置_实验案例一：交换机的初始配置

交换机的基本原理与配置_实验案例一：交换机的初始配置

1、实验环境 实验用具包括一台Cisco交换机&#xff0c;一台PC&#xff0c;一根Console 线缆。 2、需求描述 如图5.17所示&#xff0c;实验案例一的配置需求如下。 通过PC连接并配置一台Cisco交换机。在交换机的各个配置模式之间切换。将交换机主机的名称改为BDON 3、推荐步…
阅读更多...
OpenHarmony应用编译 - 如何在源码中编译复杂应用（4.0-Release）

OpenHarmony应用编译 - 如何在源码中编译复杂应用（4.0-Release）

文档环境 开发环境&#xff1a;Windows 11 编译环境&#xff1a;Ubuntu 22.04 开发板型号&#xff1a;DAYU 200&#xff08;RK3568&#xff09; 系统版本&#xff1a;OpenHarmony-4.0-Release 功能简介 在 OpenHarmony 系统中预安装应用的 hap 包会随系统编译打包到镜像中&a…
阅读更多...
C语言—每日选择题—Day68

C语言—每日选择题—Day68

第一题 1、运行以下C语言代码&#xff0c;输出的结果是&#xff08;&#xff09; #include <stdio.h> int main() {char *str[3] {"stra", "strb", "strc"};char *p str[0];int i 0;while(i < 3){printf("%s ",p);i;} retur…
阅读更多...
Path Aggregation Network for Instance Segmentation

Path Aggregation Network for Instance Segmentation

PANet 摘要1. 引言2.相关工作3.框架 PANet 最初是为 proposal-based 实例分割框架提出来的&#xff0c;mask 是实例的掩码&#xff0c;覆盖了物体包含的所有像素&#xff0c;proposal 在目标检测领域是可能存在目标的区域。在实例分割中&#xff0c;首先利用RPN(Region Proposa…
阅读更多...
【并发】第四篇 原子操作系列-AtomicInteger原子操作类详解

【并发】第四篇 原子操作系列-AtomicInteger原子操作类详解

导航 一. 简介二. 源码分析三. 原子操作原理三. 实际用途1. 标志位2. 唯一标识生成器3. 计数器一. 简介 AtomicInteger是Java中提供的一种线程安全的原子操作类，用来实现对整数类型的原子操作。它可以在多线程环境下保证对整数的原子性操作，而不需要使用synchronized关键字或…
阅读更多...
小样本计数网络FamNet(Learning To Count Everything)

小样本计数网络FamNet(Learning To Count Everything)

小样本计数网络FamNet(Learning To Count Everything) 大多数计数方法都仅仅针对一类特定的物体&#xff0c;如人群计数、汽车计数、动物计数等。一些方法可以进行多类物体的计数&#xff0c;但是training set中的类别和test set中的类别必须是相同的。 为了增加计数方法的可拓…
阅读更多...
CloudCompare——win11配置CloudComPy

CloudCompare——win11配置CloudComPy

CloudComPy配置 1 基本环境介绍2 安装Anaconda2.1 下载anaconda2.2 安装anaconda2.3 配置镜像源2.4 更改虚拟环境的默认创建位置2.5 其他问题2.5.1 激活自己创建的环境提示&#xff1a;系统找不到指定的路径2.5.2 InvalidVersionSpecError: Invalid version spec: 2.72.5.3 卸载…
阅读更多...
Hibernate框架的搭建

Hibernate框架的搭建

Hibernate框架的搭建 分层体系结构与持久化 三层体系结构 分层体系结构 指的是将系统的组件分隔到不同的层中&#xff0c;每一层中的组件应保持内聚性&#xff1b; 每一层都应与它下面的各层保持松散耦合。 层与层之间存在自上而下的依赖关系&#xff0c;即上层组件会访问下…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023