CUDA Unity Compute Shader 3

计划
这应该是第3章的读书笔记，但是因为第3章读起来比较困难，所以先看了《CUDA并行程序设计编程指南》的第5章和第6章，感觉读起来顺畅多了，《CUDA并行程序设计编程指南》暂定精读第5、6、7章

1.如何生成ptx文件
属性->CUDA C/C++ ->Common->Keep Preprocess Files->是（–keep）
在这里插入图片描述2.查看内核使用的寄存器数量；
属性->CUDA C/C++ ->Device->Verbose PTXAS Output->是 (–ptxas-options=-v)
在这里插入图片描述# NVIDIA_CUDA_Programming_Guide_1.1_chs

4.2.2 同步函数
void __syncthreas();
在一个块内同步所有线程。一旦所有线程到达了这点，恢复正常执行。
__syncthreads()通常用于调整在相同块之间的线程通信。当在一个块内的有些线程访问相同的共享或全局内存时，对于有些内存访问潜在存在read-after-write,write-after-read,或者write-after-write的危险。
这些数据危险可以通过同步线程之间的访问得以避免。
__syncthreads()允许放在条件代码中，但只有当整个线程块有相同的条件贯穿时，否则代码执行可能被挂起或导致没想到的副作用。

注解：这是对__syncthreads()函数的解释

Professional CUDA C Programming

Chapter03 CUDA Execution Model

3.3 并行性的表现

注解：和书中使用的测试程序有所不同，偷个懒，后面看时间再同步

3.3.1 用nvprof检测活跃的线程数

一个内核的可实现占用率被定义为：
每周期内活跃线程束的平均数量与一个SM支持的线程束最大数量的比值。

nvprof --metrics achieved_occupany CUDA.exe

注解：–metrics后面需要跟一个指令
在这里插入图片描述

3.3.2 用nvprof检测内存操作

gld_throughput指标检查内核的内存读取效率
gld_efficiency指标检测全局加载效率，即被请求的全局加载吞吐量占所需的全局加载吞吐量的比值。它衡量了应用程序的加载操作利用设备内存带宽的程度。

3.3.3 增大并行性

指标与性能
▨ 在大部分情况下，一个单独的指标不能产生最佳的性能
▨ 与总体性能最直接相关的指标或事件取决于内核代码的本质
▨ 在相关的指标与事件之间寻求一个好的平衡
▨ 从不同角度查看内核以寻求相关指标间的平衡
▨ 网格/块启发式算法为性能调节提供了一个很好的起点

3.4 避免分支分化

3.4.1 并行规约问题

要对一个有N个元素的整数数组求和。
▨ 相邻配对：元素与它们直接相邻的元素配对
▨ 交错配对：根据给定的跨度配对元素
C语言利用递归实现的一个交错配对方法：

int recursiveReduce(int* data, int const size) {if (size == 1)return data[0];int const stride = size / 2;for (size_t i = 0; i < stride; i++){data[i] += data[i + stride];}return recursiveReduce(data, stride);
}

在向量中执行满足交换律和结合律的运算，被称为归约问题。
3.4.2 并行归约中的分化

在这个内核里，有两个全局内存数组：一个大数组用来存放整个数组，进行归约；另一个小数组用来存放每个线程块的部分和。每个线程块在数组的一部分上独立地执行操作。
循环中迭代一次执行一个归约步骤。归约是在就地完成的，这意味着在每一步，全局内存的值都被部分和替代。
两个相邻元素间的距离被称为跨度，初始化均为。在每一次归约循环结束后，这个间隔就被乘以2。在第一次循环结束后，idata（全局数据指针）的偶数元素将被部分和替代。在第二次循环结束后，idata的每四个元素将会被新产生的部分和替代。
因为线程间无法同步，所以每个线程块产生的部分和被赋值回了主机，并且在哪儿进行串行求和。

3.4.3 改善并行归约的分化

3.4.3 交错配对的归约

3.5 展开循环

循环展开是一个尝试通过减少分支出现的频率和循环维护指令来优化循环的技术。
在循环展开中，循环主题在代码中药多次被编写，而不是只编写一次循环主题再使用另一个循环来反复执行的。
任何的封闭循环可将它的迭代次数减少或完全删除。
循环体的复制数量被称为循环展开因子，迭代次数就变为了原始循环迭代次数除以循环展开因此。

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];
}

3.5.1 展开的归约

3.5.2 展开线程的归约

3.5.3 完全展开的归约

3.5.4 模板函数的归约

#include "CUDA_Header.cuh"int invokeKernel();int recursiveReduce(int* data, int const size) {if (size == 1)return data[0];int const stride = size / 2;for (size_t i = 0; i < stride; i++){data[i] += data[i + stride];}return recursiveReduce(data, stride);
}__global__ void reduceNeighbored(int* g_idata, int* g_odata, unsigned int n) {unsigned int tid = threadIdx.x;unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;int* idata = g_idata + blockIdx.x * blockDim.x;if (idx >= n)return;for (size_t stride = 1; stride < blockDim.x; stride *= 2){int index = 2 * stride * tid;if (index < blockDim.x)idata[index] += idata[index + stride];//if ((tid % (2 * stride)) == 0)//	idata[tid] += idata[tid + stride];__syncthreads();}if (tid == 0)g_odata[blockIdx.x] = idata[0];
}__global__ void reduceInterleaved(int* g_idata, int* g_odata, unsigned int n) {unsigned int tid = threadIdx.x;unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;int* idata = g_idata + blockIdx.x * blockDim.x;if (idx >= n)return;int stride = blockDim.x / 2;for (size_t i = stride; i > 0; i >>= 1){if (tid < i)idata[tid] += idata[tid + i];__syncthreads();}if (tid == 0)g_odata[blockIdx.x] = idata[0];
}__global__ void reduceUnrolling2(int* g_idata, int* g_odata, unsigned int n) {unsigned int tid = threadIdx.x;unsigned int idx = blockIdx.x * blockDim.x * 2 + threadIdx.x;//printf("tid %d , idx %d , blockIdx.x %d, blockDim.x %d , unrolling target %d \n",tid, idx, blockIdx.x, blockDim.x, idx + blockDim.x);int* idata = g_idata + blockIdx.x * blockDim.x * 2;if (idx + blockDim.x < n)g_idata[idx] += g_idata[idx + blockDim.x];__syncthreads();//__syncthreads()语句可以保证，线程块中的任一线程在进入下一次迭代之前，在当前迭代里，每个线程的所有部分和都被保存在了全局内存中，进入下一次迭代的所有线程都使用上一步产生的数值。在最后一个循环以后，整个线程块的和被保存进全局内存中//注解：__syncthreads()的意思应该是所有线程束中的//write result for this block to global mem//注解：当迭代结束后，结果保存在了idata的第一个元素里，idata是g_idata偏移后的地址if (tid == 0)g_odata[blockIdx.x] = idata[0];
}__global__ void reduceUnrollWarp8(int* g_idata, int* g_odata, unsigned int n) {unsigned int tid = threadIdx.x;unsigned int idx = blockIdx.x * blockDim.x * 8 + threadIdx.x;int* idata = g_idata + blockIdx.x * blockDim.x * 8;//unrolling 8if (idx + 7 * blockDim.x < n) {int a1 = g_idata[idx];int a2 = g_idata[idx + blockDim.x];int a3 = g_idata[idx + 2 * blockDim.x];int a4 = g_idata[idx + 3 * blockDim.x];int b1 = g_idata[idx + 4 * blockDim.x];int b2 = g_idata[idx + 5 * blockDim.x];int b3 = g_idata[idx + 6 * blockDim.x];int b4 = g_idata[idx + 7 * blockDim.x];g_idata[idx] = a1 + a2 + a3 + a4 + b1 + b2 + b3 + b4;}__syncthreads();for (size_t stride = blockDim.x / 2; stride > 32; stride >>= 1){if (tid < stride)idata[tid] += idata[tid + stride];__syncthreads();}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];}if (tid == 0)g_odata[blockIdx.x] = idata[0];
}__global__ void reduceCompleteUnrollWarp8(int* g_idata, int* g_odata, unsigned int n) {unsigned int tid = threadIdx.x;unsigned int idx = blockIdx.x * blockDim.x * 8 + threadIdx.x;int* idata = g_idata + blockIdx.x * blockDim.x * 8;if (idx + 7 * blockDim.x < n) {int a1 = g_idata[idx];int a2 = g_idata[idx + blockDim.x];int a3 = g_idata[idx + 2 * blockDim.x];int a4 = g_idata[idx + 3 * blockDim.x];int b1 = g_idata[idx + 4 * blockDim.x];int b2 = g_idata[idx + 5 * blockDim.x];int b3 = g_idata[idx + 6 * blockDim.x];int b4 = g_idata[idx + 7 * blockDim.x];g_idata[idx] = a1 + a2 + a3 + a4 + b1 + b2 + b3 + b4;}__syncthreads();if (blockDim.x >= 1024 && tid < 512)idata[tid] += idata[tid + 512];__syncthreads();if (blockDim.x >= 512 && tid < 256)idata[tid] += idata[tid + 256];__syncthreads();if (blockDim.x >= 256 && tid < 128)idata[tid] += idata[tid + 128];__syncthreads();if (blockDim.x >= 128 && tid < 64)idata[tid] += idata[tid + 64];__syncthreads();if (tid < 32) {volatile int* vsmem = idata;vsmem[tid] += vsmem[tid + 32];vsmem[tid] += vsmem[tid + 16];vsmem[tid] += vsmem[tid + 8];vsmem[tid] += vsmem[tid + 4];vsmem[tid] += vsmem[tid + 2];vsmem[tid] += vsmem[tid + 1];}if (tid == 0)g_odata[blockIdx.x] = idata[0];
}template <unsigned int iBlockSize>
__global__ void reduceCompleteUnroll(int* g_idata, int* g_odata, unsigned int n) {unsigned int tid = threadIdx.x;unsigned int idx = blockIdx.x * blockDim.x * 8 + threadIdx.x;int* idata = g_idata + blockIdx.x * blockDim.x*8;if (idx + 7 * blockDim.x < n) {int a1 = g_idata[idx];int a2 = g_idata[idx + blockDim.x];int a3 = g_idata[idx + 2 * blockDim.x];int a4 = g_idata[idx + 3 * blockDim.x];int b1 = g_idata[idx + 4 * blockDim.x];int b2 = g_idata[idx + 5 * blockDim.x];int b3 = g_idata[idx + 6 * blockDim.x];int b4 = g_idata[idx + 7 * blockDim.x];g_idata[idx] = a1 + a2 + a3 + a4 + b1 + b2 + b3 + b4;}__syncthreads();if (blockDim.x >= 1024 && tid < 512)	idata[tid] += idata[tid + 512];		__syncthreads();if (blockDim.x >= 512 && tid < 256)		idata[tid] += idata[tid + 256];		__syncthreads();if (blockDim.x >= 256 && tid < 128)		idata[tid] += idata[tid + 128];		__syncthreads();if (blockDim.x >= 128 && tid < 64)		idata[tid] += idata[tid + 64];		__syncthreads();if (tid < 32) {volatile int* vsmem = idata;vsmem[tid] += vsmem[tid + 32];vsmem[tid] += vsmem[tid + 16];vsmem[tid] += vsmem[tid + 8];vsmem[tid] += vsmem[tid + 4];vsmem[tid] += vsmem[tid + 2];vsmem[tid] += vsmem[tid + 1];}if (tid == 0)g_odata[blockIdx.x] = idata[0];
}
int main() {invokeKernel();}static int invokeKernel() {int dev = 0;cudaDeviceProp deviceProp;CHECK(cudaGetDeviceProperties(&deviceProp, dev));printf("device %d: %s \n", dev, deviceProp.name);CHECK(cudaSetDevice(dev));bool bResult = false;long   size = 1 << 24;printf("	with array size %d \n", size);int blockSize = 1024;dim3 block(blockSize, 1);dim3 grid((size + block.x - 1) / block.x, 1);grid.x /= 8;printf("grid %d block %d \n", grid.x, block.x);size_t bytes = size * sizeof(int);int* h_idata = (int*)malloc(bytes);int* h_odata = (int*)malloc(grid.x * sizeof(int));int* tmp = (int*)malloc(bytes);for (size_t i = 0; i < size; i++){h_idata[i] = (int)(rand() & 0xFF);}memcpy(tmp, h_idata, bytes);clock_t iStart, iElaps;int gpu_sum = 0;int* d_idata = NULL;int* d_odata = NULL;CHECK(cudaMalloc((void**)&d_idata, bytes));CHECK(cudaMalloc((void**)&d_odata, grid.x * sizeof(int)));iStart = cpuSeconds();int cpu_sum = recursiveReduce(tmp, size);iElaps = cpuSeconds() - iStart;printf("cpu reduce elapsed %d sec cpu_sum: %d \n", iElaps, cpu_sum);//kernel 1: reduceNeighboredCHECK(cudaMemcpy(d_idata, h_idata, bytes, cudaMemcpyHostToDevice));CHECK(cudaDeviceSynchronize());iStart = cpuSeconds();reduceCompleteUnroll<1024> << <grid.x, block >> > (d_idata, d_odata, size);CHECK(cudaDeviceSynchronize());iElaps = cpuSeconds() - iStart;CHECK(cudaMemcpy(h_odata, d_odata, grid.x * sizeof(int), cudaMemcpyDeviceToHost));gpu_sum = 0;for (size_t i = 0; i < grid.x; i++){gpu_sum += h_odata[i];}printf("gpu Kernel elapsed %d sec gpu_sum: %d <<<grid %d, block %d>>> \n", iElaps, gpu_sum, grid.x, block.x);free(h_idata);free(h_odata);cudaFree(d_idata);cudaFree(d_odata);
}

对于以上所讲的内容，使用Excel模拟了下内核的运算过程，应该更加容易理解了
带颜色的部分并非计算结果，而是处理的数据索引
在这里插入图片描述

严重性	代码	说明	项目	文件	行	禁止显示状态
错误		kernel launch from __device__ or __global__ functions requires separate compilation mode	CUDA	H:\C_CPP_CUDA\CUDA\CUDA\nestedHelloWorld.cu	15

属性->CUDA C/C++ ->Common->Generate Relocatable Device Code-> 是（-rdc=true）
在这里插入图片描述
使用nvvp查看代码，NVIDIA Visual Profile在

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.5\bin

路径下有个nvvp.bat，双击可以打开这个窗口
从这个bat中也可以看到，nvvp.exe的路径是在

C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v12.5\libnvvp

下
在这里插入图片描述

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

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

相关文章

[QT] MAC使用Qt Creator运行程序如何仅运行一个进程？

[QT] MAC使用Qt Creator运行程序如何仅运行一个进程？

大家刚开始使用QtCreator会发现每次run程序&#xff0c;都会出现一个程序进程&#xff0c;使得调试操作增加。如下&#xff0c;每次run都会出现一个demo14的进程。 如何每次run后&#xff0c;就关闭上一次的进程&#xff0c;而重新拉起新进程呢&#xff1f; 看这里 这是默认…
阅读更多...
Redis-02

Redis-02

redis安装包位置 /opt/redis-7.2.5 redis默认安装路径&#xff1a; 配置文件路径&#xff1a;/usr/local/bin/redisconfig gcc安装位置 /opt/rhredis启动&#xff1a; 在/usr/local/bin目录下输入redis-server redisconfig/redis.confredis性能测试命令 redis-benchmark [opt…
阅读更多...
【WEEK14】 【DAY5】Swagger第三部分【中文版】

【WEEK14】 【DAY5】Swagger第三部分【中文版】

2024.5.31 Friday 接上文【WEEK14】 【DAY4】Swagger第二部分【中文版】 目录 16.6.配置API分组16.6.1.修改SwaggerConfig.java16.6.2.重启 16.7.实体配置16.7.1.新建pojo文件夹16.7.2.修改HelloController.java16.7.3.重启 16.8.常用注解16.8.1.Swagger的所有注解定义在io.swa…
阅读更多...
NIUSHOP多商户V6版预售背后的前端技术革新

NIUSHOP多商户V6版预售背后的前端技术革新

随着电子商务的快速发展&#xff0c;多商户电商平台成为了市场上的热门选择。在这个背景下&#xff0c;NIUSHOP多商户V6版的预售活动引发了广泛关注。本文将从前端技术的角度&#xff0c;探讨NIUSHOP多商户V6版在预售背后所蕴含的技术革新和亮点。 一、引言 NIUSHOP多商户系统…
阅读更多...
node.js点餐系统app-计算机毕业设计源码84406

node.js点餐系统app-计算机毕业设计源码84406

摘 要 21世纪的今天&#xff0c;随着社会的不断发展与进步&#xff0c;人们对于信息科学化的认识&#xff0c;已由低层次向高层次发展&#xff0c;由原来的感性认识向理性认识提高&#xff0c;管理工作的重要性已逐渐被人们所认识&#xff0c;科学化的管理&#xff0c;使信息存…
阅读更多...
php反序列化学习（3）

php反序列化学习（3）

1、session 当session_start()被调用或者php.ini中session.auto_start为1时&#xff0c;php内部调用会话管理器&#xff0c;访问用户session被序列化后&#xff0c;存储到指定目录&#xff08;默认为/tmp&#xff09;。 漏洞产生&#xff1a;写入格式与读取格式不一致 处理器…
阅读更多...
如何设置让背景颜色不包括 padding 部分，顺带全面学习 background-clip 属性（可以实现文字渐变）

如何设置让背景颜色不包括 padding 部分，顺带全面学习 background-clip 属性（可以实现文字渐变）

先解决需求 实现背景颜色不包括 padding 部分&#xff0c;直接给容器添加 css 属性&#xff1a;background-clip:content-box; 示例代码&#xff1a; .content-box-example {background-color: lightblue;padding: 20px;border: 1px solid black;background-clip: content-bo…
阅读更多...
06C内存分配

06C内存分配

C零碎语法 目录 文章目录 C零碎语法1.内存布局2. 内存对齐2.1结构体内存对齐2.1应用 1.内存布局 2. 内存对齐 2.1结构体内存对齐 三条原则&#xff1a; &#xff08;1&#xff09;结构体变量的 起始地址能够被其最宽的成员大小整除。 &#xff08;2&#xff09;结构体每个…
阅读更多...
前端Vue自定义滚动卡片组件设计与实现

前端Vue自定义滚动卡片组件设计与实现

摘要 随着技术的日新月异&#xff0c;前端开发的复杂度不断提升。传统的整块应用开发方式在面对小的改动或功能增加时&#xff0c;常常需要修改大量的整体逻辑&#xff0c;造成开发效率低下和维护成本高昂。为了应对这一挑战&#xff0c;组件化开发应运而生。本文将以Vue框架下…
阅读更多...
Vue3项目练习详细步骤(第二部分：主页面搭建)

Vue3项目练习详细步骤(第二部分：主页面搭建)

主页面搭建 页面主体结构 路由 子路由 主页面搭建 页面主体结构 在vuews目录下新建Layout.vue文件 主页面内容主体代码 <script setup> import {Management,Promotion,UserFilled,User,Crop,EditPen,SwitchButton,CaretBottom } from element-plus/icons-vue imp…
阅读更多...
YOLOv5改进（六）--引入YOLOv8中C2F模块

YOLOv5改进（六）--引入YOLOv8中C2F模块

文章目录 1、前言2、C3模块和C2F模块2.1、C3模块2.2、BottleNeck模块2.3、C2F模块 3、C2F代码实现3.1、common.py3.2、yolo.py3.3、yolov5s_C2F.yaml 4、目标检测系列文章 1、前言 本文主要使用YOLOv8的C2F模块替换YOLOv5中的C3模块&#xff0c;经过实验测试&#xff0c;发现Y…
阅读更多...
C++之类的三种继承修饰符（public、private、protected）总结

C++之类的三种继承修饰符（public、private、protected）总结

1、前言 前文博客介绍了修饰符public、private、protected在类中成员变量和函数访问权限限制的总结&#xff0c;主要如下&#xff1a; public&#xff08;公有&#xff09;: 公有成员在类的内部和外部都可以被访问。 private&#xff08;私有&#xff09;: 私有成员只能在类的内…
阅读更多...
Qt xml学习之calculator-qml

Qt xml学习之calculator-qml

1.功能说明&#xff1a;制作简易计算器 2.使用技术&#xff1a;qml,scxml 3.项目效果&#xff1a; 4.qml部分&#xff1a; import Calculator 1.0 //需要引用对应类的队友版本 import QtQuick 2.12 import QtQuick.Window 2.12 import QtQuick.Controls 1.4 import QtScxml…
阅读更多...
Linux 的权限

Linux 的权限

目录 Linux 的用户 root 用户 和 普通用户 如何新建普通用户&#xff1f; 如何切换用户&#xff1f; 一开始是以 root 用户登录&#xff1a; 一开始以普通用户登录&#xff1a; 如何删除用户&#xff1f; Linux文件权限 什么是读权限&#xff08; r &#xff09;&#…
阅读更多...
vs2019 c++20 规范的 STL 库的智能指针 shared、unique 、weak 、auto 及 make_** 函数的源码注释汇总，和几个结论

vs2019 c++20 规范的 STL 库的智能指针 shared、unique 、weak 、auto 及 make_** 函数的源码注释汇总，和几个结论

智能指针的源码都在 《memory》 头文件中。因为头文件太长&#xff0c;再者本次整理是基于以前的零散的模板分析。故相当于抽取了该头文件中关于智能指针的源码进行分析&#xff0c;注释。 &#xff08;1 探讨一&#xff09;当独占指针指向数组时&#xff0c;其默认的删除器是…
阅读更多...
ABP框架+Mysql（二）

ABP框架+Mysql（二）

展示页面--图书列表页面 本地化 开始的UI开发之前,我们首先要准备本地化的文本(这是你通常在开发应用程序时需要做的).本地化文本在前端页面会常用。 本地化文本位于 Acme.BookStore.Domain.Shared 项目的 Localization/BookStore 文件夹下: 打开 en.json (英文翻译)文件并更…
阅读更多...
STM32-- GPIO-＞EXTI-＞NVIC中断

STM32-- GPIO-＞EXTI-＞NVIC中断

一、NVIC简介 什么是 NVIC &#xff1f; NVIC 即嵌套向量中断控制器&#xff0c;全称 Nested vectored interrupt controller 。它 是内核的器件&#xff0c;所以它的更多描述可以看内核有关的资料。M3/M4/M7 内核都是支持 256 个中断&#xff0c;其中包含了 16 个系统中…
阅读更多...
MySQL—函数—数值函数（基础）

MySQL—函数—数值函数（基础）

一、引言 首先了解一下常见的数值函数哪些&#xff1f;并且直到它们的作用&#xff0c;并且演示这些函数的使用。 二、数值函数 常见的数值函数如下&#xff1a; 注意&#xff1a; 1、ceil(x)、floor(x) &#xff1a;向上、向下取整。 2、mod(x,y)&#xff1a;模运算&#x…
阅读更多...
CTF本地靶场搭建——GZ:CTF基础使用

CTF本地靶场搭建——GZ:CTF基础使用

GZ::CTF 是一个基于 ASP.NET Core 的开源 CTF 平台。 简介 GZ::CTF 是一个基于 ASP.NET Core 的开源 CTF 平台&#xff0c;采用 Docker 或 K8s 作为容器部署后端&#xff0c;提供了可自定义的题目类型、动态容器和动态分值功能。 本项目缘起于作者认为 CTFd 的实现不优雅&a…
阅读更多...
Photoshop 首选项设置建议

Photoshop 首选项设置建议

Windows Ps菜单&#xff1a;编辑/首选项 Edit/Preferences 快捷键&#xff1a;Ctrl K macOS Ps菜单&#xff1a;Photoshop/首选项 Photoshop/Preferences 快捷键&#xff1a;Cmd K 对 Photoshop 的首选项 Preferences进行设置&#xff0c;可以提高修图与设计效率。下面是一些…
阅读更多...
推荐文章
最新文章

Copyright @ 2022~2023