基于矩阵乘的CUDA编程优化过程

news2025/1/23 22:43:54

背景:网上很多关于矩阵乘的编程优化思路,本着看理论分析万遍,不如实际代码写一遍的想法,大概过一下优化思路。

矩阵乘的定义如下,约定矩阵的形状及存储方式为: A[M, K], B[K, N], C[M, N]。

C_{i,j}=\sum_{k=0}^{n}A_{ik}\times B_{kj}

CPU篇

朴素实现方法

        按照常规的思路,实现矩阵乘时如下的3层for循环。

#define OFFSET(row, col, ld) ((row) * (ld) + (col))
void cpuSgemm(float *a, float *b, float *c, const int M, const int N, const int K) 
{
    for (int m = 0; m < M; m++) {
        for (int n = 0; n < N; n++) {
            float psum = 0.0;
            for (int k = 0; k < K; k++) {
                psum += a[OFFSET(m, k, K)] * b[OFFSET(k, n, N)];
            }
            c[OFFSET(m, n, N)] = psum;
        }
    }
}

数据访存连续的优化

        矩阵B的存储默认为N方向连续,所以可以将上面的第2,3层循环互换顺序,这样B的取数就不会跨行了,而是连续取数,达到访问连续的效果。

void cpuSgemm_1(float *a, float *b, float *c, const int M, const int N, const int K) 
{
    for (int m = 0; m < M; m++) {
        for (int k = 0; k < K; k++) {
            for (int n = 0; n < N; n++)
            {
                c[OFFSET(m, n, N)] += a[OFFSET(m, k, K)] * b[OFFSET(k, n, N)];
            }           
        }
    }
}

数据重排/数据复用的优化

        上面将M,N,K的for循环调整为M,K,N的循环顺序,导致我们K方向累加不能缓存了,增加了多次访问C矩阵的开销,所以我们不放先直接将B矩阵转置处理,然后再按照原始的M,N,K的for循环来处理。

void cpuSgemm_2(float *a, float *b, float *c, const int M, const int N, const int K) 
{
    float* b1=(float*) malloc(sizeof(float)*K*N);
    for(int i=0; i<K; i++)
    {
        for (int j=0; j<N; j++)
        {
            b1[OFFSET(j,i,K)]= b[OFFSET(i,j,N)];
        }
    }

    for (int m = 0; m < M; m++) {
        for (int n = 0; n < N; n++) {
            float psum = 0.0;
            for (int k = 0; k < K; k++) {
                psum += a[OFFSET(m, k, K)] * b1[OFFSET(n, k, K)];
            }
            c[OFFSET(m, n, N)] = psum;
        }
    }
}

性能表现

        如下是测试CPU环境下这几种方法的时间情况,其中M=N=512, K =256。可以发现经过优化后的代码在时间上是逐步减少的。

        CPU的优化思路还有其他的,比如循环展开,intrinsic函数,基于cache的矩阵切分等,注意本文并没有都实现出来。

cpuSgemm, Time measured: 416889 microseconds.
cpuSgemm_1, Time measured: 405259 microseconds.
cpuSgemm_2, Time measured: 238786 microseconds.

GPU篇

grid线程循环矩阵乘法

        输出矩阵C有M*N个点,每个点是K个数的乘积和,所以可以定义每个线程计算K个点的乘积和,即grid线程循环矩阵乘法。

__global__ void matrix_multiply_gpu_0(float*a, float*b, float*c, int M, int N, int K)
{
    int tidx =threadIdx.x;
    int bidx = blockIdx.x;
    int idx = bidx * blockDim.x +tidx;
    int row = idx/N;
    int col = idx%N;
    if(row<M && col < N)
    {
        float tmp =0.0;
        for(int k=0; k<K; k++)
        {
            tmp+=a[row*K+k] * b[k*N+col];
        }
        c[row*N+col] = tmp;
    }
}

block线程循环矩阵乘法

        grid内线程循环的矩阵乘法有如下缺憾:一个block内线程可能需要计算C矩阵不同行的矩阵元素,block内thread对相应的A矩阵访存不一致,导致无法广播和额外的访存开销,导致执行时间增加。

        针对这个问题,可以做如下改进:每个block计算C矩阵的一行,block内的thread以固定跳步步长blockDim.x的方法循环计算C矩阵的一行,每一行启动一个block,共计M个block。

__global__ void matrix_multiply_gpu_1(float*a, float*b, float*c, int M, int N, int K)
{
    int tidx =threadIdx.x;
    int bidx = blockIdx.x;

    float tmp;
    for(;bidx<M; bidx += gridDim.x)
    {
        for(;tidx<N; tidx+=blockDim.x )
        {
            tmp=0.0;
            for(int k=0; k<K; k++)
            {
                tmp+=a[bidx*K +k] * b[k*N+tidx];
            }
            c[bidx*N+tidx] = tmp;
        }              
    }
}

行共享存储矩阵乘法

        共享存储与L1 Cache同级,其访存延迟较全局存储小一个量级。用共享存储代替全局存储是GPU最重要的优化手段之一。采用共享存储优化的关键是数据复用,数据复用次数越多,共享存储优化可获得的收益也越高。

        在block循环乘法中,1个block内所有thread都会用到A矩阵的一行,此时与B矩阵每一列相乘,A矩阵中该行复用了N次。故可以考虑将A矩阵的一行读入shared memory,运算时候从shared memory读取相应的数据。

        注意代码中TILE_WIDTH>=K。

#define TILE_WIDTH 256
__global__ void matrix_multiply_gpu_2(float*a, float*b, float*c, int M, int N, const int K)
{
    __shared__ float data[TILE_WIDTH];
    int tid = threadIdx.x;
    int row = blockIdx.x;
    int i,j;
    for(i=tid; i<K; i+=blockDim.x)
    {
        data[i]=a[row*K +i];
    }
    __syncthreads();
    float tmp;
    for(j=tid; j<N; j+=blockDim.x)
    {
        tmp=0.0;
        for(int k=0; k<K; k++)
        {
            tmp += data[k]*b[k*N+j];
        }
        c[row*N+j] = tmp;
    }
}

分块共享存储矩阵乘法

        根据上面共享存储的理解,我们很自然的想到把B矩阵也考虑数据复用,所以可以同时把A,B矩阵都分成棋盘似的小尺寸的数据块,从全局内存读取到共享内存,这样可以有效降低数据访问时间,充分复用矩阵乘的局部数据。

#define TILE_SIZE 32
__global__ void matrix_multiply_gpu_3(float*a, float*b, float*c, int M, int N, const int K)
{
    __shared__ float matA[TILE_SIZE][TILE_SIZE];
	__shared__ float matB[TILE_SIZE][TILE_SIZE];
	
	int bx = blockIdx.x;
	int by = blockIdx.y;
	int tx = threadIdx.x;
	int ty = threadIdx.y;
	
    int Col = bx * TILE_SIZE + tx;
	int Row = by * TILE_SIZE + ty;
	
	float Pervalue = 0.0;
	for(int i = 0;i < K / TILE_SIZE;i++)  
	{
		matA[ty][tx] = a[Row * K + (i * TILE_SIZE + tx)];
		matB[ty][tx] = b[Col + (i * TILE_SIZE + ty) * N];
		__syncthreads();
	
		for(int k = 0;k < TILE_SIZE;k++) 
			Pervalue += matA[ty][k] * matB[k][tx];
		__syncthreads();
	}
	
	c[Row * N + Col] = Pervalue;
    
}

性能表现

利用nvprof工具,统计各个核函数的执行时间如下,可以发现每一步优化思路都能直观的带来的性能提升。

完整代码:

GitHub - Briwisdom/study_CUDA_examples: some demos for study CUDA program.

#include <iostream>
#include <chrono>

using namespace std;

#define OFFSET(row, col, ld) ((row) * (ld) + (col))

void initDate(float *arr,int Len, bool randFlag=true)
{
    if (randFlag)
    {
        for (int i = 0; i < Len; i++) {
            arr[i] = rand()/1000000;
        }
    }
    else
    {
        float value =0.0;
        for (int i = 0; i < Len; i++) {
            arr[i] = value;
        }
    }  
}

void compare_result(float *x, float *y, int n, char *name)
{
    int cnt=0;
    for (int i=0; i<n; i++)
    {
        if (x[i]!=y[i])
        {
            cnt++;
            printf("x= %f, y= %f\n", x[i],y[i]);
        }
            
    }
    printf("%s, ", name);
    if(cnt ==0)
        printf("result matched.\n");
    else
        printf("something error! result not match number = %d int total number: %d .\n", cnt, n);

}


void cpuSgemm(float *a, float *b, float *c, const int M, const int N, const int K) 
{
    for (int m = 0; m < M; m++) {
        for (int n = 0; n < N; n++) {
            float psum = 0.0;
            for (int k = 0; k < K; k++) {
                psum += a[OFFSET(m, k, K)] * b[OFFSET(k, n, N)];
            }
            c[OFFSET(m, n, N)] = psum;
        }
    }
}

void cpuSgemm_1(float *a, float *b, float *c, const int M, const int N, const int K) 
{
    for (int m = 0; m < M; m++) {
        for (int k = 0; k < K; k++) {
            for (int n = 0; n < N; n++)
            {
                c[OFFSET(m, n, N)] += a[OFFSET(m, k, K)] * b[OFFSET(k, n, N)];
            }           
        }
    }
}

void cpuSgemm_2(float *a, float *b, float *c, const int M, const int N, const int K) 
{
    float* b1=(float*) malloc(sizeof(float)*K*N);
    for(int i=0; i<K; i++)
    {
        for (int j=0; j<N; j++)
        {
            b1[OFFSET(j,i,K)]= b[OFFSET(i,j,N)];
        }
    }

    for (int m = 0; m < M; m++) {
        for (int n = 0; n < N; n++) {
            float psum = 0.0;
            for (int k = 0; k < K; k++) {
                psum += a[OFFSET(m, k, K)] * b1[OFFSET(n, k, K)];
            }
            c[OFFSET(m, n, N)] = psum;
        }
    }
}




void operation(void (*func)(float*,float*, float*, int, int, int), float *a, float *b, float *c, const int M, const int N, const int K, int repeat, char* name)
{
    auto begin0 = std::chrono::high_resolution_clock::now();
    for(int i=0; i<repeat; i++)
    {
        (*func)(a,b,c, M, N, K);
    }
    auto end0 = std::chrono::high_resolution_clock::now();
    auto elapsed0 = std::chrono::duration_cast<std::chrono::microseconds>(end0 - begin0);
    printf("%s, Time measured: %d microseconds.\n", name, int(elapsed0.count()/repeat));
}

__global__ void matrix_multiply_gpu_0(float*a, float*b, float*c, int M, int N, int K)
{
    int tidx =threadIdx.x;
    int bidx = blockIdx.x;
    int idx = bidx * blockDim.x +tidx;
    int row = idx/N;
    int col = idx%N;
    if(row<M && col < N)
    {
        float tmp =0.0;
        for(int k=0; k<K; k++)
        {
            tmp+=a[row*K+k] * b[k*N+col];
        }
        c[row*N+col] = tmp;
    }
}

__global__ void matrix_multiply_gpu_1(float*a, float*b, float*c, int M, int N, int K)
{
    int tidx =threadIdx.x;
    int bidx = blockIdx.x;

    float tmp;
    for(;bidx<M; bidx += gridDim.x)
    {
        for(;tidx<N; tidx+=blockDim.x )
        {
            tmp=0.0;
            for(int k=0; k<K; k++)
            {
                tmp+=a[bidx*K +k] * b[k*N+tidx];
            }
            c[bidx*N+tidx] = tmp;
        }              
    }
}

#define TILE_WIDTH 256
__global__ void matrix_multiply_gpu_2(float*a, float*b, float*c, int M, int N, const int K)
{
    __shared__ float data[TILE_WIDTH];
    int tid = threadIdx.x;
    int row = blockIdx.x;
    int i,j;
    for(i=tid; i<K; i+=blockDim.x)
    {
        data[i]=a[row*K +i];
    }
    __syncthreads();
    float tmp;
    for(j=tid; j<N; j+=blockDim.x)
    {
        tmp=0.0;
        for(int k=0; k<K; k++)
        {
            tmp += data[k]*b[k*N+j];
        }
        c[row*N+j] = tmp;
    }
}

#define TILE_SIZE 32
__global__ void matrix_multiply_gpu_3(float*a, float*b, float*c, int M, int N, const int K)
{
    __shared__ float matA[TILE_SIZE][TILE_SIZE];
	__shared__ float matB[TILE_SIZE][TILE_SIZE];
	
	int bx = blockIdx.x;
	int by = blockIdx.y;
	int tx = threadIdx.x;
	int ty = threadIdx.y;
	
    int Col = bx * TILE_SIZE + tx;
	int Row = by * TILE_SIZE + ty;
	
	float Pervalue = 0.0;
	for(int i = 0;i < K / TILE_SIZE;i++)  
	{
		matA[ty][tx] = a[Row * K + (i * TILE_SIZE + tx)];
		matB[ty][tx] = b[Col + (i * TILE_SIZE + ty) * N];
		__syncthreads();
	
		for(int k = 0;k < TILE_SIZE;k++) 
			Pervalue += matA[ty][k] * matB[k][tx];
		__syncthreads();
	}
	
	c[Row * N + Col] = Pervalue;
    
}
 

int main()
{
    int M=512;
    int N=512;
    int K=256;

    float *a = (float*) malloc(M*K * sizeof(float));
    float *b = (float*) malloc(N*K * sizeof(float));
    float *c = (float*) malloc(M*N * sizeof(float));
    float *c1 = (float*) malloc(M*N * sizeof(float));
    float *c2 = (float*) malloc(M*N * sizeof(float));
    float *c_gpu_0 = (float*) malloc(M*N * sizeof(float));
    float *c_gpu_1 = (float*) malloc(M*N * sizeof(float));
    float *c_gpu_2 = (float*) malloc(M*N * sizeof(float));
    float *c_gpu_3 = (float*) malloc(M*N * sizeof(float));

    initDate(a,M*K);
    initDate(b,N*K);
    initDate(c, M*N, false);
    initDate(c1, M*N, false);
    initDate(c2, M*N, false);
    initDate(c_gpu_0, M*N, false);
    initDate(c_gpu_1, M*N, false);
    initDate(c_gpu_2, M*N, false);
    initDate(c_gpu_3, M*N, false);

    //ensure result is right.
    cpuSgemm(a,b,c,M,N,K);
    cpuSgemm_1(a,b,c1,M,N,K);
    cpuSgemm_2(a,b,c2,M,N,K); 
    compare_result(c, c1, M*N,"sgemm1");
    compare_result(c, c2,  M*N,"sgemm2");



    //test the prerformance.
    int repeat =10;
    operation(cpuSgemm,a,b,c,M,N,K,repeat,"cpuSgemm");
    operation(cpuSgemm_1,a,b,c1,M,N,K,repeat,"cpuSgemm_1");
    operation(cpuSgemm_2,a,b,c2,M,N,K,repeat,"cpuSgemm_2");
    
    float* d_a, *d_b, *d_c0, *d_c1, *d_c2, *d_c3;
    cudaMalloc((void**) &d_a, sizeof(float)*(M*K));
    cudaMalloc((void**) &d_b, sizeof(float)*(N*K));
    cudaMalloc((void**) &d_c0, sizeof(float)*(M*N));
    cudaMalloc((void**) &d_c1, sizeof(float)*(M*N));
    cudaMalloc((void**) &d_c2, sizeof(float)*(M*N));
    cudaMalloc((void**) &d_c3, sizeof(float)*(M*N));

    cudaMemcpy(d_a, a, sizeof(float)*M*K, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, sizeof(float)*N*K, cudaMemcpyHostToDevice);
    

    int threadnum=64;
    int blocks =(M*N+threadnum-1)/threadnum;
    cudaMemcpy(d_c0, c_gpu_0, sizeof(float)*M*N, cudaMemcpyHostToDevice);
    matrix_multiply_gpu_0<<<blocks, threadnum>>>(d_a, d_b, d_c0, M, N, K);
    cudaMemcpy(c_gpu_0, d_c0, sizeof(float)*M*N, cudaMemcpyDeviceToHost);
    compare_result(c, c_gpu_0,  M*N,"gpu_0");
    cudaFree(d_c0);

    cudaMemcpy(d_c1, c_gpu_1, sizeof(float)*M*N, cudaMemcpyHostToDevice);
    matrix_multiply_gpu_1<<<M, threadnum>>>(d_a, d_b, d_c1, M, N, K);
    cudaMemcpy(c_gpu_1, d_c1, sizeof(float)*M*N, cudaMemcpyDeviceToHost);
    compare_result(c, c_gpu_1,  M*N,"gpu_1");
    cudaFree(d_c1);

    cudaMemcpy(d_c2, c_gpu_2, sizeof(float)*M*N, cudaMemcpyHostToDevice);
    matrix_multiply_gpu_2<<<M, threadnum>>>(d_a, d_b, d_c2, M, N, K);
    cudaMemcpy(c_gpu_2, d_c2, sizeof(float)*M*N, cudaMemcpyDeviceToHost);
    compare_result(c, c_gpu_2,  M*N,"gpu_2");
    cudaFree(d_c2);

    threadnum=32;
    dim3 gridSize(M / threadnum,N / threadnum);
	dim3 blockSize(threadnum,threadnum);
    cudaMemcpy(d_c3, c_gpu_3, sizeof(float)*M*N, cudaMemcpyHostToDevice);
    matrix_multiply_gpu_3<<<gridSize, blockSize>>>(d_a, d_b, d_c3, M, N, K);
    cudaMemcpy(c_gpu_3, d_c3, sizeof(float)*M*N, cudaMemcpyDeviceToHost);
    compare_result(c, c_gpu_3,  M*N,"gpu_3");
    cudaFree(d_c3);


    free(a);
    free(b);
    free(c);
    free(c1);
    free(c2);
    free(c_gpu_0);
    free(c_gpu_1);
    free(c_gpu_2);
    free(c_gpu_3);
    cudaFree(d_a);
    cudaFree(d_b);

}

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

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

相关文章

阿里云域名外部入库流程

注册商是阿里云&#xff0c;且在阿里云管理的&#xff0c;请使用此教程外部入库。 如您的域名注册商是阿里云但在聚名管理&#xff0c;请参考教程&#xff1a;https://www.west.cn/faq/list.asp?unid2539 在外部入库操作之前&#xff0c;请先登录阿里云获取账号ID。详细的账…

SPSS25软件安装包下载及安装教程

SPSS 25下载链接&#xff1a;https://docs.qq.com/doc/DUlFFZ0dpWVhUZFpW 1.选中下载好的安装包&#xff0c;鼠标右键解压到“SPSS 25 64bit”文件夹 2.选中”SPPS 25 64bit.exe“鼠标右击选择以管理员身份运行 3.点击“下一步” 4.选择“我接受许可协议中的全部条款”&#x…

Hive06_基础查询

HIVE 查询语句 1 查询语句语法&#xff1a; SELECT [ALL | DISTINCT] select_expr, select_expr, ... FROM table_reference [WHERE where_condition] [GROUP BY col_list] [ORDER BY col_list] [CLUSTER BY col_list | [DISTRIBUTE BY col_list] [SORT BY col_list] ] [LIMI…

YOLOv8改进:IoU系列篇 | Shape-IoU关注边界框本身的形状和尺度来计算损失 | 2023年12月最新IoU改进

🚀🚀🚀本文改进: 提出了一种新颖的Shape-IoU,小目标检测实现涨点,更加关注边界框本身的形状和尺度来计算损失 🚀🚀🚀YOLOv8改进专栏:http://t.csdnimg.cn/hGhVK 学姐带你学习YOLOv8,从入门到创新,轻轻松松搞定科研; 1.Shape-IoU原理介绍 论文:https://ar…

window下载安装Mongodb数据库

我们先要访问他的官网 https://www.mongodb.com/zh-cn 然后顶部导航栏 选择 (Products/产品) 下的 (Community Edition/社区版) 进入界面后 找到 MongoDB Community Server Download 点击下面的按钮 Select package 然后会弹到这个位置 第一个版本 用系统默认选择的就好 第二…

第三百四十一回

文章目录 1. 概念介绍2. 使用方法与主要功能2.1 使用方法2.2 主要功能 3. 示例代码4. 内容总结 我们在上一章回中介绍了"如何获取App自身信息"相关的内容&#xff0c;本章回中将介绍一个三方包:open_setting.闲话休提&#xff0c;让我们一起Talk Flutter吧。 1. 概念…

mysql报错:can‘t create more than max_prepared_stmt_count statements

max_prepared_stmt_count 参数控制了一个 MySQL 实例能够准备的最大预处理语句&#xff08;prepared statements&#xff09;的数量。 预处理语句是一种优化技术&#xff0c;可以在应用程序发送sql语句到数据库之前先将其编译和缓存起来&#xff0c;以提高sql的执行效率以及防…

【Amazon Bedrock】体验 Bedrock 的基本功能,为构建强大安全的LLM应用而准备

文章目录 一、什么是Amazon Bedrock&#xff1f;二、为什么选择 Amazon Bedrock三、访问Amazon Bedrock UI四、与Amazon Bedrock 聊天五、对比Amazon Bedrock 不同基础模型的返回结果六、让Amazon Bedrock处理文本七、利用Amazon Bedrock生成图片八、参考链接 一、什么是Amazon…

CMake入门教程【基础篇】打印(message)

文章目录 1. 基本用法示例 2. 打印变量的值示例 3. 打印列表的值示例 4. 打印生成器表达式的值示例 5.总结 #mermaid-svg-pXC2tr41PvHonKJa {font-family:"trebuchet ms",verdana,arial,sans-serif;font-size:16px;fill:#333;}#mermaid-svg-pXC2tr41PvHonKJa .error-…

解决Golang WriteHeader设置后,Content-Type失效的问题

场景 最近笔者在研究web框架过程中&#xff0c;发现了一个响应类型的问题&#xff0c;困扰许久&#xff0c;原因就是设置了响应状态码后&#xff0c;然后设置响应类型为application/json。在实际请求后&#xff0c;响应类型变成了text/plain; charsetutf-8格式。 问题解决&…

redis安装与配置(Ubuntu)

目录 1. 切换到 root 用户 2. 搜索安装包 3. 安装 redis 4. 查看 redis 是否正常存在 5. 修改ip 6. 重新启动服务器 7. 连接服务器 1. 切换到 root 用户 通过 su 命令切换到 root 用户。 2. 搜索安装包 apt search redis 这里安装的是下面的版本&#xff1a; 3. 安装 …

三、C语言中的分支与循环—while循环 (5)

本章分支结构的学习内容如下&#xff1a; 三、C语言中的分支与循环—if语句 (1) 三、C语言中的分支与循环—关系操作符 (2) 三、C语言中的分支与循环—条件操作符 与逻辑操作符(3) 三、C语言中的分支与循环—switch语句&#xff08;4&#xff09;分支结构 完 本章循环结…

2024.1.1力扣每日一题——经营摩天轮的最大利润

2024.1.1 题目来源我的题解方法一 模拟 题目来源 力扣每日一题&#xff1b;题序&#xff1a;1599 我的题解 方法一 模拟 计算当前上摩天轮的人数和等待的人数就可以得到该轮次的利润&#xff0c;然后一只更新最大利润就可以了。 时间复杂度&#xff1a;O(n)。n数组的长度 空…

关于LayUI表格重载数据问题

目的 搜索框搜索内容重载数据只显示搜索到的结果 遇到的问题 在layui官方文档里介绍的table属性有data项,但使用下列代码 table.reload(test, {data:data //data为json数据}); 时发现&#xff0c;会会重新调用table.render的url拿到原来的数据&#xff0c;并不会显示出来传…

DDoS 攻击并不是全部来自于PC组成的僵尸网络

DDoS&#xff0c;分布式拒绝服务攻击&#xff0c;是指处于不同位置的多个攻击者同时向一个或数个目标发动攻击&#xff0c;或者一个攻击者控制了位于不同位置的多台机器并利用这些机器对受害者同时实施攻击。很多人会以为DDoS 攻击&#xff0c;全都是攻击者控制PC肉鸡发起的攻击…

Android Studio xml布局代码补全功能失效问题

这里写目录标题 前言&#xff1a;问题描述原因分析&#xff1a;解决方案&#xff1a;1.更新 Android Studio 版本2.原版本解决XML补全失效 小结 前言&#xff1a; 在开发过程中&#xff0c;你可能遇到很多奇奇怪怪的问题。Android Studio 编译器出现问题也是常有的事情&#x…

2 Windows网络编程

1 基础概念 1.1 socket概念 Socket 的原意是“插座”&#xff0c;在计算机通信领域&#xff0c;socket 被翻译为“套接字”&#xff0c;它是计算机之间进行通信的一种约定或一种方式。Socket本质上是一个抽象层&#xff0c;它是一组用于网络通信的API&#xff0c;包括了一系列…

图片处理相关网站(图片分辨率、尺寸修改、AI扩图等)

分享一些免费的可进行图片的各种处理的网站&#xff0c;包括图片分辨率、尺寸修改、AI扩图等&#xff0c;持续增加中。。。 1.photokit.com 可进行图片分辨率、尺寸、压缩等修改。 免费在线图片编辑器 - 在线抠图、改图、修图、美图 - PhotoKit.comPhotoKit是一款免费的…

浅析PCIe 6.0功能更新与实现的挑战-2

确保TX重试缓冲区的准确性也非常重要&#xff0c;因为在接收到确认或否定信号之前&#xff0c;所有FLIT都需要存储在缓冲区中。由于一个FLIT可能包含多个TLP&#xff0c;或者一个大TLP可以被分割成多个FLIT&#xff0c;因此必须保证重传的FLIT不会跳过或额外添加原始FLIT中的TL…

【算法与数据结构】968、LeetCode监控二叉树

文章目录 一、题目二、解法三、完整代码 所有的LeetCode题解索引&#xff0c;可以看这篇文章——【算法和数据结构】LeetCode题解。 一、题目 二、解法 思路分析&#xff1a;本题的一共有两个难点&#xff0c;一个在于如何遍历二叉树&#xff08;前中后遍历&#xff0c;选择什么…