【CUDA】Branch Divergence and Unrolling Loop
目录
一、避免分支发散
1.1 并行规约问题
1.2 并行规约中的发散
二、UNrolling Loops
一、避免分支发散
控制流有时依赖于 thread 索引。同一个warp中,一个条件分支可能导致性能很差。通过重新组织数据获取模式可以减少或避免 warp divergence。具体问题查看下面这篇文章:
【CUDA】Warp解析-CSDN博客https://blog.csdn.net/GG_Bruse/article/details/143772619
1.1 并行规约问题
若要计算一个数组N个元素的和,使用CPU编程实现十分容易
int sum = 0;
for (int i = 0; i < N; ++i)sum += array[i]
若数组中的元素非常多,应用并行计算可以大大提高效率。鉴于加法交换律等性质,这个求和过程可以以元素的任意顺序来进行:
- 将输入数组切割成很多小的块
- 用 thread 来计算每个块的和
- 对这些块的结果再求和得最终结果
数组的切割主旨是:用 thread 求数组中按一定规律配对的的两个元素和,然后将所有结果组合成一个新的数组,然后再次求配对两元素和,多次迭代,直到数组中只有一个结果
比较直观的两种实现方式是:
- Neighbored pair:每次迭代都是相邻两个元素求和
- Interleaved pair:按一定跨度配对两个元素
对于N个元素的数组,该过程需N - 1次求和、步。InterleavedPair的跨度是半个数组长度
上述讲的这类问题术语为 reduction problem。Parallel reduction(并行规约)是指迭代减少操作,是并行算法中非常关键的一种操作
1.2 并行规约中的发散
Neighbored pair
在这个kernel中,有两个global memory array,一个用来存放数组所有数据,另一个用来存放部分和。所有 block 独立的执行求和操作
#include <iostream>
#include <cuda_runtime.h>
#include <cuda_runtime_api.h>
using namespace std;// 1 << 28 == 268435456
__global__
void reduceNeighbored(int* inputData, int *outputData, unsigned int N)
{unsigned int threadIndex = threadIdx.x; // 0 - 511unsigned int index = blockIdx.x * blockDim.x + threadIdx.x; // 0 - 524287, 512, 0 - 511 即 0 - 268435455int* iData = inputData + blockIdx.x * blockDim.x; // inputData + (0 - 268434944) 即指向每个block的起始位置if (index >= N) return;// stride = (1 - 511) 1, 2, 4, ..., 500for (int stride = 1; stride < blockDim.x; stride *= 2) {if ((threadIndex % (2 * stride)) == 0) // 每次取threadIndex为偶数的threadiData[threadIndex] += iData[threadIndex + stride];__syncthreads();}if (threadIndex == 0) outputData[blockIdx.x] = iData[0]; // 放入计算出的数据
}void GPU_ReduceNeighbored(int* dInputData, int* dOutputData, int size, dim3 grid, dim3 block)
{reduceNeighbored<<<grid, block>>>(dInputData, dOutputData, size);
}
因为无法让所有的block同步,所以最后将所有block的结果送回 host 进行串行计算
主函数代码:
void GPU_ReduceNeighbored(int* inputData, int *outputData, int N, dim3 grid, dim3 block);long long Seconds()
{// 获取当前时间点auto now = std::chrono::high_resolution_clock::now();// 将时间点转换为以毫秒为单位的时间间隔auto duration = std::chrono::duration_cast<std::chrono::milliseconds>(now.time_since_epoch());// 获取毫秒数return duration.count();
}int main(int argc, char** argv)
{// set up deviceint device = 0;cudaDeviceProp deviceProp;cudaGetDeviceProperties(&deviceProp, device);printf("%s starting reduction at device %d : %s\n", argv[0], device, deviceProp.name);cudaSetDevice(device);// initialization int size = 1 << 28;printf("Array Size : %d\n", size);int blockSize = 512;dim3 block(blockSize, 1);dim3 grid((size + block.x - 1) / block.x, 1);printf("grid %d , block %d\n", grid.x, block.x);// allocate host memorysize_t bytes = size * sizeof(int);int* hInputData = (int*) malloc(bytes);int* hOutputData = (int*) malloc(grid.x * sizeof(int));int* tmp = (int*) malloc(bytes);for(int i = 0; i < size; ++i) hInputData[i] = (int)(rand() & 0xFF); // 屏蔽最大两个字节memcpy(tmp, hInputData, bytes);// allocate device memoryint* dInputData = nullptr;int* dOutputData = nullptr;cudaMalloc((void **) &dInputData, bytes);cudaMalloc((void **) &dOutputData, grid.x * sizeof(int));int iStart, iElaps;// kernel: reduceNeighboredcudaMemcpy(dInputData, hInputData, bytes, cudaMemcpyHostToDevice);cudaDeviceSynchronize();iStart = Seconds();GPU_ReduceNeighbored(dInputData, dOutputData, size, grid, block);cudaDeviceSynchronize();iElaps = Seconds() - iStart;cudaMemcpy(hOutputData, dOutputData, grid.x * sizeof(int), cudaMemcpyDeviceToHost);long long gpuSum = 0;for (int i = 0; i < grid.x; i++) gpuSum += hOutputData[i];printf("gpu Neighbored elapsed %d ms gpuSum: %lld <<<grid %d block %d>>>\n", iElaps, gpuSum, grid.x, block.x);// free host memoryfree(hInputData);free(hOutputData);// free device memorycudaFree(dInputData);cudaFree(dOutputData);// reset devicecudaDeviceReset();return 0;
}
if ((tid % (2 * stride)) == 0) 该表达式只对偶数ID的线程为true,所以其导致很高的 divergent warps。第一次迭代只有偶数ID的线程执行了指令,但是所有线程都要被调度;第二次迭代,只有四分之一的 thread 是 active 的,但所有 thread 都要被调度。可以重新组织每个线程对应的数组索引来强制ID相邻的thread来处理求和操作
__global__
void reduceNeighboredLess(int* inputData, int* outputData, unsigned int N)
{unsigned int threadIndex = threadIdx.x; // 0 - 511unsigned int index = blockIdx.x * blockDim.x + threadIdx.x; // 0 - 524287, 512, 0 - 511 即 0 - 268435455int* iData = inputData + blockIdx.x * blockDim.x; // inputData + (0 - 268434944) 即指向每个block的起始位置if(index >= N) return;// stride = (1 - 511) 1, 2, 4, ..., 500for (int stride = 1; stride < blockDim.x; stride *= 2) {int idx = 2 * stride * threadIndex; // 2 * (1 - 500) * (0 - 511)if (idx < blockDim.x) iData[idx] += iData[idx + stride];__syncthreads(); }if (threadIndex == 0) outputData[blockIdx.x] = iData[0];
}void GPU_ReduceNeighboredLess(int* dInputData, int* dOutputData, int size, dim3 grid, dim3 block)
{reduceNeighboredLess<<<grid, block>>>(dInputData, dOutputData, size);
}
int index = 2 * stride * tid; 因为步调乘以了2,下面的语句使用block的前半部分thread来执行求和:if (index < blockDim.x)
对于一个有512个 thread 的 block 而言,前八个 warp 执行第一轮 reduction,剩下八个 warp 什么也不干;第二轮,前四个 warp 执行,剩下十二个什么也不干。就不存在divergence了(divergence只发生于同一个warp)。但最后还是会有divergence,因为这个时候需要执行 threads 已经凑不够一个 warp 了
Interleaved pair
Interleaved Pair模式的初始步调是block大小的一半,每个thread处理相隔半个block的两个数据
__global__
void reduceInterleaved(int *inputData, int *outputData, unsigned int N)
{unsigned int threadIndex = threadIdx.x;unsigned int index = blockIdx.x * blockDim.x + threadIdx.x;int *iData = inputData + blockIdx.x * blockDim.x;if(index >= N) return;for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {if (threadIndex < stride)iData[threadIndex] += iData[threadIndex + stride];__syncthreads();}if (threadIndex == 0) outputData[blockIdx.x] = iData[0];
}void GPU_ReduceInterleaved(int* dInputData, int* dOutputData, int size, dim3 grid, dim3 block)
{reduceInterleaved<<<grid, block>>>(dInputData, dOutputData, size);
}
步调被初始化为block大小的一半:for (int stride = blockDim.x / 2; stride > 0; stride >>= 1)
下面的语句使得第一次迭代时,block的前半部分thread执行相加操作,第二次是前四分之一,以此类推:if (tid < stride)
二、UNrolling Loops
loop unrolling(取消循环)是用来优化循环减少分支的方法,该方法就是把本应在多次 loop 中完成的操作,尽量压缩到一次 loop。循环体展开程度称为loop unrolling factor(循环展开因子),loop unrolling对顺序数组的循环操作性能有很大影响
for (int i = 0; i < 100; ++i) {a[i] = b[i] + c[i];
}
如下重复一次循环体操作,迭代数目将减少一半:
for (int i = 0; i < 100; i += 2) {a[i] = b[i] + c[i];a[i+1] = b[i+1] + c[i+1];
}
从高级语言层面是无法看出性能提升的原因的,需从 low-level instruction 层面去分析,第二段代码循环次数减少了一半,而循环体两句语句的读写操作的执行在CPU上是可以同时执行互相独立的,所以相对第一段,第二段性能要好
Unrolling 在CUDA编程中意义更重。目标是通过减少指令执行消耗,增加更多的独立指令来提高性能。这样就会增加更多的并行操作从而产生更高的指令和内存带宽(bandwidth)。也就提供了更多的 eligible warps 来帮助 hide instruction / memory latency
在前文的 reduceInterleaved 中,每个block处理一部分数据,将这些数据称为 dataBlock
下面的代码是 reduceInterleaved 的修正版本,每个 block 都是以两个 dataBlock 作为源数据进行操作(前文中,每个 block 处理一个 dataBlock)。这是一种循环分区:每个 thread 作用于多个 dataBlock,并且从每个 dataBlock 中取出一个元素处理
__global__
void reduceUnrolling(int *inputData, int *outputData, unsigned int N)
{unsigned int threadIndex = threadIdx.x;unsigned int index = blockIdx.x * blockDim.x * 2 + threadIdx.x;int *iData = inputData + blockIdx.x * blockDim.x * 2;// unrolling 2 data blocksif (index + blockDim.x < N) inputData[index] += inputData[index + blockDim.x];__syncthreads();for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {if (threadIndex < stride) iData[threadIndex] += iData[threadIndex + stride];__syncthreads();}if (threadIndex == 0) outputData[blockIdx.x] = iData[0];
} void GPU_ReduceUnrolling(int* dInputData, int* dOutputData, int size, dim3 grid, dim3 block)
{reduceUnrolling<<<grid.x / 2, block>>>(dInputData, dOutputData, size);
}
注意下面的语句,每个 thread 从相邻的 dataBlock中取数据,这一步实际上就是将两个 dataBlock规约成一个:if (idx + blockDim.x < n) g_idata[idx] += g_idata[idx+blockDim.x];
global array index 也要相应的调整,因为,相对之前的版本,同样的数据,只需要原来一半的 thread 就能解决问题。要注意的是,这样做也会降低warp或block的并行性(因为thread减少):
由于每个block处理两个data block,所以需要调整grid的配置:
reduceUnrolling2<<<grid.x / 2, block>>>(d_idata, d_odata, size);
同一个 thread 中若能有更多的独立的 load/store 操作,会产生更好的性能,因为这样做 memory latency 能够更好的被隐藏
__syncthreads是用来同步block内部thread的。在reduction kernel中,其被用来在每次循环中保证所有 thread 的写 global memory 的操作都已完成,这样才进行下一阶段的计算
当kernel进行到只需要少于或等32个thread(即一个warp)呢?由于使用的SIMT模式,warp内的thread 是有一个隐式的同步过程的。最后几次迭代可以用下面的语句展开:
if (tid < 32) {volatile int *vmem = idata;vmem[tid] += vmem[tid + 32];vmem[tid] += vmem[tid + 16];vmem[tid] += vmem[tid + 8];vmem[tid] += vmem[tid + 4];vmem[tid] += vmem[tid + 2];vmem[tid] += vmem[tid + 1];
}
warp unrolling避免了__syncthreads同步操作,因为这一步本身就没必要
注意volatile修饰符,其告诉编译器每次执行赋值时必须将 vmem[tid] 的值 store 回 global memory。若不这样,编译器或 cache 可能会优化读写global/shared memory。有了这个修饰符,编译器就会认为这个值会被其他 thread 修改,从而使得每次读写都直接去 memory 而不是去 cache 或者 register
__global__ void reduceUnrollWarps8(int *inputData, int *outputData, unsigned int N)
{unsigned int threadIndex = threadIdx.x;unsigned int index = blockIdx.x * blockDim.x * 8 + threadIdx.x;int *iData = inputData + blockIdx.x * blockDim.x * 8;// unrolling 8if (index + 7*blockDim.x < N) {int a1 = inputData[index];int a2 = inputData[index + blockDim.x];int a3 = inputData[index + 2 * blockDim.x];int a4 = inputData[index + 3 * blockDim.x];int b1 = inputData[index + 4 * blockDim.x];int b2 = inputData[index + 5 * blockDim.x];int b3 = inputData[index + 6 * blockDim.x];int b4 = inputData[index + 7 * blockDim.x];inputData[index] = a1 + a2 + a3 + a4 + b1 + b2 + b3 + b4;}__syncthreads();for (int stride = blockDim.x / 2; stride > 32; stride >>= 1) {if (threadIndex < stride) iData[threadIndex] += iData[threadIndex + stride];__syncthreads();}// unrolling warpif (threadIndex < 32) {volatile int *vmem = iData;vmem[threadIndex] += vmem[threadIndex + 32];vmem[threadIndex] += vmem[threadIndex + 16];vmem[threadIndex] += vmem[threadIndex + 8];vmem[threadIndex] += vmem[threadIndex + 4];vmem[threadIndex] += vmem[threadIndex + 2];vmem[threadIndex] += vmem[threadIndex + 1];}if (threadIndex == 0) outputData[blockIdx.x] = iData[0];
} void GPU_ReduceUnrolling8(int* dInputData, int* dOutputData, int size, dim3 grid, dim3 block)
{reduceUnrollWarps8<<<grid.x / 8, block>>>(dInputData, dOutputData, size);
}
若在编译时已知了迭代次数,就可以完全把循环展开。Fermi 和 Kepler 每个 block 的最大 thread 数目都是1024,文中的 kernel 的迭代次数都是基于 blockDim 的,所以完全展开循环是可行的
__global__
void reduceCompleteUnrollWarps8(int *inputData, int *outputData, unsigned int N)
{unsigned int threadIndex = threadIdx.x;unsigned int index = blockIdx.x * blockDim.x * 8 + threadIdx.x;int *iData = inputData + blockIdx.x * blockDim.x * 8;// unrolling 8if (index + 7 * blockDim.x < N) {int a1 = inputData[index];int a2 = inputData[index + blockDim.x];int a3 = inputData[index + 2 * blockDim.x];int a4 = inputData[index + 3 * blockDim.x];int b1 = inputData[index + 4 * blockDim.x];int b2 = inputData[index + 5 * blockDim.x];int b3 = inputData[index + 6 * blockDim.x];int b4 = inputData[index + 7 * blockDim.x];inputData[index] = a1 + a2 + a3 + a4 + b1 + b2 + b3 + b4;}__syncthreads();if (blockDim.x>=1024 && threadIndex < 512) iData[threadIndex] += iData[threadIndex + 512];__syncthreads(); if (blockDim.x>=512 && threadIndex < 256) iData[threadIndex] += iData[threadIndex + 256];__syncthreads();if (blockDim.x>=256 && threadIndex < 128) iData[threadIndex] += iData[threadIndex + 128];__syncthreads();if (blockDim.x>=128 && threadIndex < 64) iData[threadIndex] += iData[threadIndex + 64];__syncthreads();// unrolling warpif (threadIndex < 32) {volatile int *vsmem = iData;vsmem[threadIndex] += vsmem[threadIndex + 32];vsmem[threadIndex] += vsmem[threadIndex + 16];vsmem[threadIndex] += vsmem[threadIndex + 8];vsmem[threadIndex] += vsmem[threadIndex + 4];vsmem[threadIndex] += vsmem[threadIndex + 2];vsmem[threadIndex] += vsmem[threadIndex + 1];}if (threadIndex == 0) outputData[blockIdx.x] = iData[0];
} void GPU_ReduceCompleteUnrollWarps8(int* dInputData, int* dOutputData, int size, dim3 grid, dim3 block)
{reduceCompleteUnrollWarps8<<<grid.x / 8, block>>>(dInputData, dOutputData, size);
}
相关文章:

【CUDA】Branch Divergence and Unrolling Loop
目录 一、避免分支发散 1.1 并行规约问题 1.2 并行规约中的发散 二、UNrolling Loops 一、避免分支发散 控制流有时依赖于 thread 索引。同一个warp中,一个条件分支可能导致性能很差。通过重新组织数据获取模式可以减少或避免 warp divergence。具体问题查看下…...
深度学习:卷积神经网络的计算复杂度,顺序操作,最大路径长度
卷积层的计算复杂度 在深度学习中,卷积层的计算复杂度主要取决于卷积核的大小、输入和输出的通道数量、以及输入序列的长度。具体来说,卷积层的计算复杂度可以通过以下几个因素来计算: 卷积核大小 k:卷积核的大小决定了每次卷积操…...
springboot 配置文件中 multipart.max-file-size 各个版本的写法
由于springboot具有几个版本,不同版本对于文件上传最大限制的配置也有所不同。 所以要注意springboot本身的版本,不然会一直报错 在springboot1.3版本中: multipart.maxFileSize在springboot1.4与springboot1.5版本中: spring…...

linux 中mysql查看慢日志
1、到mysql容器,先登录到数据库,查看是否开启 mysql -h 127.0.0.1 -uroot -p SHOW VARIABLES LIKE slow_query_log; 2、如果没有开启,需要先开启 set global slow_query_log ON; 3、查看慢日志文件 SHOW VARIABLES LIKE slow_query_log…...
单片机的基本组成与工作原理
单片机(Microcontroller Unit, MCU)是一种将计算机的主要部分集成在一个芯片上的小型计算机系统。它通常包括中央处理器(CPU)、存储器(Memory)、输入输出接口(I/O Ports)、定时器/计…...
智慧隧道和智慧交通
通过引入先进的物联网技术,将各种硬件设备如传感器、摄像头、控制系统等有效地连接并管理起来,以实现道路安全和交通流畅的目标。这些设备将能够实时监控和控制隧道内的各种设备和系统,从而提高道路安全、提升驾驶体验并降低管理成本。 在这个…...
List、Set、Map详解和区别
在 Java 中,List、Set、Map是常用的集合类型,它们各自具有不同的特点和用途,以下是对它们的详细介绍及区别分析: List(列表) 特点: 有序性:List中的元素是有序的,即元素…...

界面控件DevExpress WinForms v24.2新功能预览 - 支持.NET 9
DevExpress WinForms 拥有180组件和UI库,能为Windows Forms平台创建具有影响力的业务解决方案。DevExpress WinForms能完美构建流畅、美观且易于使用的应用程序,无论是Office风格的界面,还是分析处理大批量的业务数据,它都能轻松胜…...

Postman之pm.test断言操作
Postman之pm.test断言操作 1.断言方法2.连接符3.条件判断符 用于验证请求的响应数据是否符合预期 1.断言方法 pm.test():定义一个测试函数,接受两个参数,一个字符串参数用来描述该测试,一个返回True/False的函数 语法格式&#…...
对数几率回归
对数几率回归简介 对数几率回归(Logistic Regression)是一种用于解决分类问题的经典统计模型,其核心思想是利用逻辑函数(Sigmoid函数)将线性回归模型的输出值映射到概率范围 [0, 1],从而实现分类预测。对数…...

docker 配置同宿主机共同网段的IP 同时通过通网段的另一个电脑实现远程连接docker
docker配置网络 #宿主机执行命令 ifconfig 查询对应的主机ip 子网掩码 网关地址 #[网卡名称]:inet[主机IP] netmask[子网掩码] broadcast[网关地址]这里需要重点关注:eno1[网卡名称]以及【192.168.31.225】网关地址 在宿主机执行docker命令创建一个虚拟…...
4-7-1.C# 数据容器 - LinkedList(LinkedList 的定义、LinkedList 结点的遍历、LinkedList 的常用方法)
LinkedList 概述 LinkedList<T> 通过节点(Node)来存储数据,每个节点包含数据和指向下一个节点的引用 LinkedList<T> 存储的元素是可重复的 LinkedList<T> 支持泛型,可以指定存储的元素的类型 LinkedList<…...

「三」体验HarmonyOS端云一体化开发模板——使用DevEco Studio直接创建端云一体化工程
关于作者 白晓明 宁夏图尔科技有限公司董事长兼CEO、坚果派联合创始人 华为HDE、润和软件HiHope社区专家、鸿蒙KOL、仓颉KOL 华为开发者学堂/51CTO学堂/CSDN学堂认证讲师 开放原子开源基金会2023开源贡献之星 「目录」 「一」HarmonyOS端云一体化概要 「二」体验HarmonyOS端云一…...

确保以管理员权限运行 Visual Studio 开发者命令提示符
文章目录 解决方法:1. 以管理员身份运行命令提示符2. 改变目录权限3. 改变项目目录位置4. 检查文件系统权限 总结: ********************************************************************** ** Visual Studio 2022 Developer Command Prompt v17.12.0 …...

命令执行简单(棱角社区有毒)
前言:小迪安全2022第一节反弹shell,小迪用的是两台都是云服务器,没有服务器可以在自己的主机上搭建也是可以的,主机上搭两个网站 思路:生成一个木马文件,下载到本机,然后利用本机上传到目标主机…...

Keil基于ARM Compiler 5的工程迁移为ARM Compiler 6的工程
环境: keil版本为5.38,版本务必高于5.30 STM32F4的pack包版本要高于2.9 软件包下载地址:https://zhuanlan.zhihu.com/p/262507061 一、更改Keil中编译器 更改后编译,会报很多错,先不管。 二、更改头文件依赖 观察…...
Kafka-创建topic源码
一、命令创建topic kafka-topics --create --topic quickstart-events --bootstrap-server cdh1:9092 --partitions 2 --replication-factor 2 二、kafka-topics脚本 exec $(dirname $0)/kafka-run-class.sh org.apache.kafka.tools.TopicCommand "$" 脚本中指定了…...

【网络安全】(一) 0成本添加访问级监控
互联网的安全感这个概念源于阿里。顾名思义,让互联网的用户对于web产品能够产生足够的信任和依赖。特别是涉及到用户资金交易的站点,一次严重的用户资料泄露就可以彻底毁掉你的品牌。 然而当前阶段除了bat大部分互联网行业的企业对于网络安全给的重视都…...

【Three.js基础学习】26. Animated galaxy
前言 shaders实现星系 课程回顾 使用顶点着色器为每个粒子设置动画 a属性 , u制服 ,v变化 像素比:window.devicePixelRatio 自动从渲染器检索像素比 renderer.getPixelRatio() 如何尺寸衰减, 放大缩小视角时,粒子都是同…...
vscode使用ssh配置docker容器环境
1 创建容器,并映射主机和容器的指定ssh服务端口 2 进入容器 docker exec -it <容器ID> /bin/bash 3在容器中安装ssh服务 apt-get update apt-get install openssh-server 接着修改ssh文件信息,将容器的10008端口暴露出来允许root用户使用ssh登录 vim /…...

【人工智能】神经网络的优化器optimizer(二):Adagrad自适应学习率优化器
一.自适应梯度算法Adagrad概述 Adagrad(Adaptive Gradient Algorithm)是一种自适应学习率的优化算法,由Duchi等人在2011年提出。其核心思想是针对不同参数自动调整学习率,适合处理稀疏数据和不同参数梯度差异较大的场景。Adagrad通…...
管理学院权限管理系统开发总结
文章目录 🎓 管理学院权限管理系统开发总结 - 现代化Web应用实践之路📝 项目概述🏗️ 技术架构设计后端技术栈前端技术栈 💡 核心功能特性1. 用户管理模块2. 权限管理系统3. 统计报表功能4. 用户体验优化 🗄️ 数据库设…...
基于Java Swing的电子通讯录设计与实现:附系统托盘功能代码详解
JAVASQL电子通讯录带系统托盘 一、系统概述 本电子通讯录系统采用Java Swing开发桌面应用,结合SQLite数据库实现联系人管理功能,并集成系统托盘功能提升用户体验。系统支持联系人的增删改查、分组管理、搜索过滤等功能,同时可以最小化到系统…...

AI病理诊断七剑下天山,医疗未来触手可及
一、病理诊断困局:刀尖上的医学艺术 1.1 金标准背后的隐痛 病理诊断被誉为"诊断的诊断",医生需通过显微镜观察组织切片,在细胞迷宫中捕捉癌变信号。某省病理质控报告显示,基层医院误诊率达12%-15%,专家会诊…...
现有的 Redis 分布式锁库(如 Redisson)提供了哪些便利?
现有的 Redis 分布式锁库(如 Redisson)相比于开发者自己基于 Redis 命令(如 SETNX, EXPIRE, DEL)手动实现分布式锁,提供了巨大的便利性和健壮性。主要体现在以下几个方面: 原子性保证 (Atomicity)ÿ…...
LRU 缓存机制详解与实现(Java版) + 力扣解决
📌 LRU 缓存机制详解与实现(Java版) 一、📖 问题背景 在日常开发中,我们经常会使用 缓存(Cache) 来提升性能。但由于内存有限,缓存不可能无限增长,于是需要策略决定&am…...

基于PHP的连锁酒店管理系统
有需要请加文章底部Q哦 可远程调试 基于PHP的连锁酒店管理系统 一 介绍 连锁酒店管理系统基于原生PHP开发,数据库mysql,前端bootstrap。系统角色分为用户和管理员。 技术栈 phpmysqlbootstrapphpstudyvscode 二 功能 用户 1 注册/登录/注销 2 个人中…...

零知开源——STM32F103RBT6驱动 ICM20948 九轴传感器及 vofa + 上位机可视化教程
STM32F1 本教程使用零知标准板(STM32F103RBT6)通过I2C驱动ICM20948九轴传感器,实现姿态解算,并通过串口将数据实时发送至VOFA上位机进行3D可视化。代码基于开源库修改优化,适合嵌入式及物联网开发者。在基础驱动上新增…...
【FTP】ftp文件传输会丢包吗?批量几百个文件传输,有一些文件没有传输完整,如何解决?
FTP(File Transfer Protocol)本身是一个基于 TCP 的协议,理论上不会丢包。但 FTP 文件传输过程中仍可能出现文件不完整、丢失或损坏的情况,主要原因包括: ✅ 一、FTP传输可能“丢包”或文件不完整的原因 原因描述网络…...

车载诊断架构 --- ZEVonUDS(J1979-3)简介第一篇
我是穿拖鞋的汉子,魔都中坚持长期主义的汽车电子工程师。 老规矩,分享一段喜欢的文字,避免自己成为高知识低文化的工程师: 做到欲望极简,了解自己的真实欲望,不受外在潮流的影响,不盲从,不跟风。把自己的精力全部用在自己。一是去掉多余,凡事找规律,基础是诚信;二是…...