【CUDA】 扫描 Scan

Scan

Scan操作是许多应用程序中常见的操作。扫描操作采用一个二元运算符⊕和一个输入数组并计算输出数组如下：

[x0,(x0⊕x1),…,( x0⊕x1⊕…..⊕xn-1)]

分层扫描和多种Scan算法介绍

Kogge-Stone's Algorithm

Kogge-Stone's Algorithm最初是为设计快速加法电路而发明的。该算法如今被用于设计高速计算机算术硬件。

图1是该算法的示例。

图1

Kogge-Stone's Algorithm，所有线程将迭代直到log2N步骤。

Kogge-Stone's Algorithm完成Scan操作所需的工作量接近Nlog2N。

Kogge-Stone's Algorithm将需要大约Nlog2N/P的时间步骤才能完成(P 执行单元的个数)。

Brent-Kung's Algorithm

Brent-Kung's Algorithm是Kogge-Stone's Algorithm的一个变种，它在计算中间结果时采取策略，以减少算法执行的工作量。

算法分为两个阶段：

降维树阶段

分发树阶段

图2是算法示例。

图2

Brent-Kung's Algorithm中，所有线程将迭代2log2N的步数。

Brent-Kung's Algorithm完成Scan操作所需的工作量接近于 2N - 2 - log2N。

Brent-Kung's Algorithm将需要大约(N / 2) * (2 * log2N - 1) / P的时间步数来完成(P 执行单元的个数)。 （由于 CUDA 设备设计该算法非常接近 Kogge-Stone's Algorithm）

三阶段Scan算法

三阶段Scan算法旨在实现比其他两种更高的工作效率。

该算法分为三个阶段：

独立扫描阶段（按块）

对先前独立扫描的最后元素进行扫描操作

将前一阶段的每个元素重新添加回块的最后阶段

图3是算法展示：

图3

具有T个线程完成的工作量是

1阶段N – 1

2阶段 Tlog2T

3阶段N- T

工作总量是 N – 1 + Tlog2T +N- T

如果我们使用P个执行单元，可以预计执行所需的时间。

花费（N – 1 + Tlog2T +N- T）/ P

分层Scan

上述描述的kernel在数据的输入大小方面有限制。

Kogge-Stone's Algorithm可以处理最多 threads-per-block个元素。

Brent-Kung's Algorithm可以处理最多2 * threads-per-block个元素。

三阶段扫描可以处理至多shared-memory / sizeof(T)个元素。

分层Scan算法旨在处理任意长度的输入。

图4是该算法展示：

图4

算法首先使用上述三种算法之一扫描输入的块，并将每个扫描的最后一个元素写入辅助数组。

然后扫描辅助数组。

并将每个元素i添加到index为 i + 1 的块中。

如果辅助数组足够大，它将使用相同的算法对其进行扫描。

Code

Host

host代码使用随机值初始化输入向量，并调用kernel执行Scan运算。

#include <iostream>
#include <cstdio>
#include <ctime>
#include <cmath>
#include <cuda_runtime.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>#include "helper_cuda.h"
#include "error.cuh"
#include "scanKernels.cuh"const double EPSILON = 1.0e-10;
const int FORTIME = 50;void check_result(thrust::host_vector<int>& h_res, thrust::host_vector<int> d_res, int n, std::string name) {bool success = true;for (int i = 0; i < n; i++) {if (abs(h_res[i] - d_res[i]) > 0.001) {std::cout << "Error at " << i << ":" << h_res[i] << "!=" << d_res[i] << std::endl;success = false;break;}}std::cout << (success ? "Success!" : "Failed..") << "(" << name << ")" << std::endl;
}int main(void)
{int n1, n2, n3, n4, n5, n6;n1 = 1024;n2 = 1024;n3 = 1024;n4 = 2048;n5 = 2048;n6 = 2048;thrust::host_vector<int> h_vec1(n1), h_vec2(n2), h_vec3(n3),h_vec4(n4), h_vec5(n5), h_vec6(n6);thrust::host_vector<int> h_res1(n1), h_res2(n2), h_res3(n3),h_res4(n4), h_res5(n5), h_res6(n6);thrust::device_vector<int> d_vec1(n1), d_vec2(n2), d_vec3(n3),d_vec4(n4), d_vec5(n5), d_vec6(n6);thrust::device_vector<int> d_res1(n1), d_res2(n2), d_res3(n3),d_res4(n4), d_res5(n5), d_res6(n6);std::fill(h_vec1.begin(), h_vec1.end(), 1);std::fill(h_vec2.begin(), h_vec2.end(), 1);std::fill(h_vec3.begin(), h_vec3.end(), 1);std::fill(h_vec4.begin(), h_vec4.end(), 1);std::fill(h_vec5.begin(), h_vec5.end(), 1);std::fill(h_vec6.begin(), h_vec6.end(), 1);thrust::inclusive_scan(h_vec1.begin(), h_vec1.end(), h_res1.begin());thrust::inclusive_scan(h_vec2.begin(), h_vec2.end(), h_res2.begin());thrust::inclusive_scan(h_vec3.begin(), h_vec3.end(), h_res3.begin());thrust::inclusive_scan(h_vec4.begin(), h_vec4.end(), h_res4.begin());thrust::inclusive_scan(h_vec5.begin(), h_vec5.end(), h_res5.begin());thrust::inclusive_scan(h_vec6.begin(), h_vec6.end(), h_res6.begin());cudaEvent_t start, stop;checkCudaErrors(cudaEventCreate(&start));checkCudaErrors(cudaEventCreate(&stop));d_vec1 = h_vec1;checkCudaErrors(cudaEventRecord(start));for (int i = 0; i < FORTIME; i++)Kogge_Stone_inclusive_scan<int> <<<1, 1024, 2048 * sizeof(int) >>> (thrust::raw_pointer_cast(d_vec1.data()), thrust::raw_pointer_cast(d_res1.data()), n1);checkCudaErrors(cudaEventRecord(stop));checkCudaErrors(cudaEventSynchronize(stop));float elapsed_time;checkCudaErrors(cudaEventElapsedTime(&elapsed_time, start, stop));printf("Time = %g ms.\n", elapsed_time / FORTIME);check_result(h_res1, d_res1, n1, std::string("Kogge Stone"));d_vec2 = h_vec1;checkCudaErrors(cudaEventRecord(start));for (int i = 0; i < FORTIME; i++)Brent_Kung_inclusive_scan<int> <<<1, 1024, 2048 * sizeof(int) >>> (thrust::raw_pointer_cast(d_vec2.data()), thrust::raw_pointer_cast(d_res2.data()), n2);checkCudaErrors(cudaEventRecord(stop));checkCudaErrors(cudaEventSynchronize(stop));checkCudaErrors(cudaEventElapsedTime(&elapsed_time, start, stop));printf("Time = %g ms.\n", elapsed_time / FORTIME);check_result(h_res2, d_res2, n2, std::string("Brent Kung"));d_vec3 = h_vec3;checkCudaErrors(cudaEventRecord(start));for (int i = 0; i < FORTIME; i++)three_phase_parallel_inclusive_scan<int> <<<1, 512, 4096 * sizeof(int) >>> (thrust::raw_pointer_cast(d_vec3.data()), thrust::raw_pointer_cast(d_res3.data()), n3, 4096);checkCudaErrors(cudaEventRecord(stop));checkCudaErrors(cudaEventSynchronize(stop));checkCudaErrors(cudaEventElapsedTime(&elapsed_time, start, stop));printf("Time = %g ms.\n", elapsed_time / FORTIME);check_result(h_res3, d_res3, n3, std::string("Three Phase Parallel"));return 0;
}

Note：

helper_cuda.h 与error.cuh头文件为错误查验工具。后续会发布到Github。

去除checkCudaErrors等错误查验函数不影响程序运行。

以下是kernel的展示。

Kogge-Stone's Scan

// This kernel can handle up to min(shared_mem_size, max threads per block) / 2 elements
template<typename T> __global__
void Kogge_Stone_inclusive_scan(T *X, T *Y, int n, T *S = NULL){extern __shared__ uint8_t shared_mem[];T *XY = reinterpret_cast<T*>(shared_mem);int i = blockIdx.x * blockDim.x + threadIdx.x;// Load elements into shared memoryif(i < n) XY[threadIdx.x] = X[i];__syncthreads();// Perform scan operationfor (unsigned int stride = 1; stride < blockDim.x; stride <<= 1){if (threadIdx.x >= stride) XY[blockDim.x + threadIdx.x] = XY[threadIdx.x] + XY[threadIdx.x - stride];__syncthreads();if (threadIdx.x >= stride) XY[threadIdx.x] = XY[blockDim.x + threadIdx.x];__syncthreads();}// Write results// Write result to global memoryif (i < n) Y[i] = XY[threadIdx.x];// Write last element to global memoryif (S != NULL && threadIdx.x == blockDim.x - 1)S[blockIdx.x] = XY[blockDim.x  - 1];
}

kernel首先将输入数组中的元素加载到共享内存中。

int i = blockIdx.x * blockDim.x + threadIdx.x;// Load elements into shared memory
if(i < n) XY[threadIdx.x] = X[i];__syncthreads();

然后它执行扫描操作。

// Perform scan operation
for (unsigned int stride = 1; stride < blockDim.x; stride <<= 1){if (threadIdx.x >= stride) XY[blockDim.x + threadIdx.x] = XY[threadIdx.x] + XY[threadIdx.x - stride];__syncthreads();if (threadIdx.x >= stride) XY[threadIdx.x] = XY[blockDim.x + threadIdx.x];__syncthreads();
}

元素写入输出数组，最后一个元素写入辅助数组。

// Write results to global memory
if (i < n) Y[i] = XY[threadIdx.x];// Write last element to global memory
if (S != NULL && threadIdx.x == blockDim.x - 1)S[blockIdx.x] = XY[blockDim.x  - 1];

Brent-Kung's Scan

// This kernel can handle up to min(shared_mem_size, max threads per block) elements
template<typename T> __global__
void Brent_Kung_inclusive_scan(T *X, T *Y, int n, T *S = NULL){extern __shared__ uint8_t shared_mem[];T *XY = reinterpret_cast<T*>(shared_mem);int i = 2 * blockIdx.x * blockDim.x + threadIdx.x;int tx = threadIdx.x;// Load elements into shared memoryXY[tx] = (i < n) ? X[i] : 0;XY[tx + blockDim.x] = (i + blockDim.x < n) ? X[i + blockDim.x] : 0;__syncthreads();// Perform scan operation (Phase 1 - Reduction)for (unsigned int stride = 1; stride <= blockDim.x; stride <<= 1){unsigned int index = (tx + 1) * stride * 2 - 1;if(index < 2 * blockDim.x) XY[index] += XY[index - stride];__syncthreads();}// Perform scan operation (Phase 2 - Reverse-Tree)for (unsigned int stride = blockDim.x / 2; stride > 0; stride >>= 1){unsigned int index = (tx + 1) * stride * 2 - 1;if(index + stride < 2 * blockDim.x) XY[index + stride] += XY[index];__syncthreads();}// Write results// Write result to global memoryif(i < n) Y[i] = XY[tx];if(i + blockDim.x < n) Y[i + blockDim.x] = XY[tx + blockDim.x];// Write last element to global memoryif (S != NULL && tx == blockDim.x - 1)S[blockIdx.x] = XY[2 * blockDim.x  - 1];
}

kernel首先从全局内存中加载两组元素。

int i = 2 * blockIdx.x * blockDim.x + threadIdx.x;
int tx = threadIdx.x;// Load elements into shared memory
XY[tx] = (i < n) ? X[i] : 0;XY[tx + blockDim.x] = (i + blockDim.x < n) ? X[i + blockDim.x] : 0;__syncthreads();

然后它进入了第一阶段的reduction。

// Perform scan operation (Phase 1 - Reduction)
for (unsigned int stride = 1; stride <= blockDim.x; stride <<= 1){unsigned int index = (tx + 1) * stride * 2 - 1;if(index < 2 * blockDim.x) XY[2 * blockDim.x + index] = XY[index] + XY[index - stride];__syncthreads();if(index < 2 * blockDim.x) XY[index] = XY[2 * blockDim.x + index];__syncthreads();
}

然后它进入了分发树的第二阶段。

// Perform scan operation (Phase 2 - Reverse-Tree)
for (unsigned int stride = blockDim.x / 2; stride > 0; stride >>= 1){unsigned int index = (tx + 1) * stride * 2 - 1;if(index + stride < 2 * blockDim.x) XY[2 * blockDim.x + index + stride] = XY[index + stride] + XY[index];__syncthreads();if(index + stride < 2 * blockDim.x) XY[index + stride] = XY[2 * blockDim.x + index + stride];__syncthreads();
}

元素被写入输出数组，最后一个元素写入辅助数组。

// Write results to global memory
if(i < n) Y[i] = XY[tx];
if(i + blockDim.x < n) Y[i + blockDim.x] = XY[tx + blockDim.x];// Write last element to global memory
if (S != NULL && tx == blockDim.x - 1)S[blockIdx.x] = XY[2 * blockDim.x  - 1];

Three-Phase-Scan Algorithm

template<typename T> __global__
void three_phase_parallel_inclusive_scan(T *X, T *Y, int n, int shared_mem_size, T *S = NULL){extern __shared__ uint8_t shared_mem[];T *XY = reinterpret_cast<T*>(shared_mem);int start_pos = blockIdx.x * shared_mem_size;int tx = threadIdx.x;for(int j = tx; j < shared_mem_size; j += blockDim.x) XY[j] = start_pos + j < n ? X[start_pos + j] : 0;__syncthreads();// Phase 1// Scan each sectionint size = shared_mem_size / blockDim.x;int start = size * tx;int stop = start + size;T acc = 0;if(start_pos + start < n) { // Threads that all their values are 0 do not execute (start > n)for(int i = start; i < stop; ++i) {acc += XY[i];XY[i] = acc;}}__syncthreads();// Phase 2// Use Brent-Kung algorithm to scan the last elements of each sectionfor (unsigned int stride = 1; stride <= blockDim.x; stride <<= 1){__syncthreads();unsigned int index = (tx + 1) * stride * 2 * size - 1;if(index < shared_mem_size) XY[index] += XY[index - (stride * size)];}for (unsigned int stride = shared_mem_size / 4; stride >= size; stride >>= 1){__syncthreads();unsigned int index = (tx + 1) * stride * 2 - 1;if(index + stride < shared_mem_size) XY[index + stride] += XY[index];}__syncthreads();// Phase 3// Add the last elements of each section to the next sectionif(tx != 0) {int value = XY[start - 1];for(int i = start; i < stop - 1; ++i) XY[i] += value;}__syncthreads();// Write results// Write result to global memoryfor(int i = tx; i < shared_mem_size; i += blockDim.x) if(start_pos + i < n) Y[start_pos + i] = XY[i];// Write last element to global memoryif (S != NULL && tx == blockDim.x - 1)S[blockIdx.x] = XY[shared_mem_size - 1];
}

kernel首先从全局内存加载元素。

int start_pos = blockIdx.x * shared_mem_size;
int tx = threadIdx.x;for(int j = tx; j < shared_mem_size; j += blockDim.x) XY[j] = start_pos + j < n ? X[start_pos + j] : 0;
__syncthreads();

然后每个线程在自己的部分执行顺序扫描。

// Phase 1
// Scan each section
int size = shared_mem_size / blockDim.x;
int start = size * tx;
int stop = start + size;
T acc = 0;
if(start_pos + start < n) { // Threads that all their values are 0 do not execute (start > n)for(int i = start; i < stop; ++i) {acc += XY[i];XY[i] = acc;}
}
__syncthreads();

然后使用Brent-Kung algorithm扫描每个部分的最后一个元素。

然后除了第一个线程外，所有线程将对应于前一个线程的元素写入它们的元素。

// Phase 3
// Add the last elements of each section to the next section
if(tx != 0) {int value = XY[start - 1];for(int i = start; i < stop - 1; ++i) XY[i] += value;
}
__syncthreads();

元素被写入输出数组，最后一个元素写入辅助数组。

// Write results// Write result to global memory
for(int i = tx; i < shared_mem_size; i += blockDim.x) if(start_pos + i < n) Y[start_pos + i] = XY[i];// Write last element to global memory
if (S != NULL && tx == blockDim.x - 1)S[blockIdx.x] = XY[shared_mem_size - 1];

分层Scan

template<typename T>
void hierarchical_scan_using_algorithm(T *X, T *Y, int n, int t_x, int shared_mem_size, int elems_per_block){T *d_S;int blocks = (n + elems_per_block - 1) / elems_per_block;// Allocate memory on the device for the intermediate resultscudaMalloc((void **)&d_S, blocks * sizeof(T));// Perform scan operationalgorithm<T><<<blocks, t_x, shared_mem_size * sizeof(T), launch.get_stream()>>>(X, Y, n, shared_mem_size, d_S);if(blocks > elems_per_block) hierarchical_scan_three_phase<T>(d_S, d_S, blocks, t_x, shared_mem_size, elems_per_block, launch);if(blocks > 1){if(blocks <= elems_per_block)algorithm<T><<<1, t_x, shared_mem_size * sizeof(T), launch.get_stream()>>>(d_S, d_S, blocks, shared_mem_size, NULL);// Add the intermediate results to the final resultsadd_intermediate_results<T><<<blocks - 1, t_x, 0, launch.get_stream()>>>(Y + elems_per_block, n - elems_per_block, d_S, elems_per_block);}// Free memory on the devicecudaFree(d_S);
}

kernel首先计算输入将被分割成多少块。

int blocks = (n + elems_per_block - 1) / elems_per_block;

然后它在设备上为中间结果分配内存。

// Allocate memory on the device for the intermediate results
cudaMalloc((void **)&d_S, blocks * sizeof(T));

然后，它对每个块执行扫描操作，并将最后一个元素写入辅助数组。

// Perform scan operation
algorithm<T><<<blocks, t_x, shared_mem_size * sizeof(T), launch.get_stream()>>>(X, Y, n, shared_mem_size, d_S);

如果块的数量大于每个块的元素数量，则递归调用该算法来扫描整个辅助数组。

if(blocks > elems_per_block) hierarchical_scan_three_phase<T>(d_S, d_S, blocks, t_x, shared_mem_size, elems_per_block, launch);

或者我们使用基本算法来扫描辅助阵列。

if(blocks <= elems_per_block)algorithm<T><<<1, t_x, shared_mem_size * sizeof(T), launch.get_stream()>>>(d_S, d_S, blocks, shared_mem_size, NULL);

最后，如果块的数量大于1(意味着我们需要分发结果)，我们调用kernel，将辅助数组添加到最终输出中。

// Add the intermediate results to the final results
add_intermediate_results<T><<<blocks - 1, t_x, 0, launch.get_stream()>>>(Y + elems_per_block, n - elems_per_block, d_S, elems_per_block);

性能分析

运行时间：

向量维度：1024

线程块维度：1024  （Three-Phase-Scan Algorithm为512）

Kogge_Stone 比 Brent_Kung 加法次数多，但是在计算资源充足情况下，纵向计算要小于Brent_Kung。计算资源充足情况下，在计算时间上Kogge_Stone 可能比 Brent_Kung 占优势，这也可能是因为Brent_Kung算法的复杂性增加。Three Phase Scan Algorithm优势在于可以计算shared-memory / sizeof(T)个元素，在不分层的情况下，一次能计算元素最多。

笔者采用设备：RTX3060 6GB

参考文献：

1、大规模并行处理器编程实战（第2版）

2、PPMP

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

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

相关文章

JavaEE——计算机工作原理

JavaEE——计算机工作原理

冯诺依曼体系&#xff08;VonNeumannArchitecture&#xff09; 现代计算机&#xff0c;大多遵守冯诺依曼体系结构 CPU中央处理器&#xff1a;进行算术运算与逻辑判断 存储器&#xff1a;分为外存和内存&#xff0c;用于存储数据&#xff08;使用二进制存储&#xff09; 输入…
阅读更多...
第一天(点亮led灯+led灯闪烁)——Arduino uno R3 学习之旅

第一天(点亮led灯+led灯闪烁)——Arduino uno R3 学习之旅

​ 常识: 一般智能手机的额定工作电流大约为200mA Arduino Uno板上I/0(输入/输出)引脚最大输出电流为40 mA Uno板控制器总的输出电流为200 mA 点亮LED灯 发光二极管介绍 发光二极管(Light Emitting Diode&#xff0c;简称LED)是一种能够将电能转化为光能的固态的半导体器件…
阅读更多...
实现模型贴图的移动缩放旋转

实现模型贴图的移动缩放旋转

技术&#xff1a;threejscanvasfabric 效果图&#xff1a; 原理&#xff1a;threejs中没有局部贴图的效果&#xff0c;只能通过map 的方式贴到模型上&#xff0c;所以说换一种方式来实现&#xff0c;通过canvasfabric来实现图片的移动缩放旋转&#xff0c;然后将整个画布以map…
阅读更多...
【STM32】在标准库中使用DMA

【STM32】在标准库中使用DMA

1.MDA简介 DMA全称Direct Memory Access,直接存储区访问。 DMA传输将数据从一个地址空间复制到另一个地址空间。当CPU初始化这个传输动作&#xff0c;传输动作本身是由DMA控制器来实现和完成的。DMA传输方式无需CPU直接控制传输&#xff0c;也没有中断处理方式那样保留现场和…
阅读更多...
【踩坑】探究PyTorch中创建稀疏矩阵的内存占用过大的问题

【踩坑】探究PyTorch中创建稀疏矩阵的内存占用过大的问题

转载请注明出处&#xff1a;小锋学长生活大爆炸[xfxuezhagn.cn] 如果本文帮助到了你&#xff0c;欢迎[点赞、收藏、关注]哦~ 目录 问题复现 原因分析 解决方案 碎碎念 问题复现 创建一个COO格式的稀疏矩阵&#xff0c;根据计算公式&#xff0c;他应该只占用约5120MB的内存&…
阅读更多...
go zero入门

go zero入门

一、goctl安装 goctl 是 go-zero 的内置脚手架&#xff0c;可以一键生成代码、文档、部署 k8s yaml、dockerfile 等。 # Go 1.16 及以后版本 go install github.com/zeromicro/go-zero/tools/goctllatest检查是否安装成功 $ goctl -v goctl version 1.6.6 darwin/amd64vscod…
阅读更多...
通过SDK使用百度智能云的图像生成模型SDXL

通过SDK使用百度智能云的图像生成模型SDXL

登录进入百度智能云控制台&#xff0c;在模型广场按照图像生成类别进行筛选&#xff0c;可以找到Stable-Diffusion-XL模型。点击Stable-Diffusion-XL模型的API文档后在弹出的新页面下拉可以找到SDK调用的说明。 import qianfandef sdxl(file: str, prompt: str, steps: int 2…
阅读更多...
C语言_练习题

C语言_练习题

求最小公倍数 思路&#xff1a;假设两个数&#xff0c;5和7&#xff0c;那么最小至少也要7吧&#xff0c;所以先假定最小公倍数是两个数之间较大的&#xff0c;然后看7能不能同时整除5和7&#xff0c;不能就加1继续除 int GetLCM(int _num1, int _num2) {int max _num1>_n…
阅读更多...
堆叠的作用

堆叠的作用

一、为什么要堆叠 传统的园区网络采用设备和链路冗余来保证高可靠性&#xff0c;但其链路利用率低、网络维护成本高&#xff0c;堆叠技术将多台交换机虚拟成一台交换机&#xff0c;达到简化网络部署和降低网络维护工作量的目的。 二、堆叠优势 1、提高可靠性 堆叠系统多台成…
阅读更多...
25款404网页源码（下）

25款404网页源码（下）

25款404网页源码&#xff08;下&#xff09; 13部分源码 14部分源码 15部分源码 16部分源码 17部分源码 18部分源码 19部分源码 20部分源码 21部分源码 22部分源码 23部分源码 24部分源码 25部分源码 领取完整源码下期更新 13 部分源码 .rail {position: absolute;width: 100%…
阅读更多...
Node.js-path 模块

Node.js-path 模块

path 模块 path 模块提供了 操作路径 的功能&#xff0c;如下是几个较为常用的几个 API&#xff1a; 代码实例&#xff1a; const path require(path);//获取路径分隔符 console.log(path.sep);//拼接绝对路径 console.log(path.resolve(__dirname, test));//解析路径 let pa…
阅读更多...
Docker搭建MySQL双主复制详细教程

Docker搭建MySQL双主复制详细教程

在此之前需要提前安装好Docker和 Docker Compose 。 一、创建目录 首先创建一个本地数据挂载目录。 mkdir -p master1-data master2-data二、编写docker-compose.yml version: 3.7services:mysql-master1:image: mysql:5.7.36container_name: mysql-master1environment:MYSQL_…
阅读更多...
mac｜idea导入通义灵码插件

mac｜idea导入通义灵码插件

官方教程&#xff1a;通义灵码下载安装指南_智能编码助手_AI编程_云效(Apsara Devops)-阿里云帮助中心 下载插件&#xff1a; ⇩ TONGYI Lingma - JetBrains 结果如下&#xff1a; 选择apply、ok&#xff0c;会出现弹窗&#xff0c;点击登录 可以实现&#xff1a;生成单元测…
阅读更多...
FRP反向隧道代理打CFS三层

FRP反向隧道代理打CFS三层

目录 攻击机 查看服务端frps.ini配置文件 开启服务端frps 蚁剑打目标机 上传客户端frp到目标机 ​frpc.ini文件配置成 客户端打开代理frpc vps显示成功客户端frpc打开 访问成功192.168.22.22的第二层内网主机 省去前面漏洞利用的rce过程&#xff0c;直接蚁剑开搞隧道…
阅读更多...
如何使用VScode创建和上传Arduino项目

如何使用VScode创建和上传Arduino项目

Visual Studio Code &#xff08;VS Code&#xff09; 是一种非常流行的通用集成开发环境 &#xff08;IDE&#xff09;。IDE 是一种将文本编辑器、编程界面、调试视图和项目管理集成在一个地方的软件。这个开源项目由微软领导&#xff0c;可以在所有操作系统上运行。使 VS Cod…
阅读更多...
深度解析Ubuntu版本升级：LTS版本升级指南

深度解析Ubuntu版本升级：LTS版本升级指南

深度解析Ubuntu版本升级&#xff1a;Ubuntu版本生命周期及LTS版本升级指南 Ubuntu是全球最受欢迎的Linux发行版之一&#xff0c;其版本升级与维护策略直接影响了无数用户的开发和生产环境。Canonical公司为Ubuntu制定了明确的生命周期和发布节奏&#xff0c;使得社区、企业和开…
阅读更多...
宿舍报修小程序的设计

宿舍报修小程序的设计

管理员账户功能包括&#xff1a;系统首页&#xff0c;个人中心&#xff0c;管理员管理&#xff0c;基础数据管理&#xff0c;论坛管理&#xff0c;故障上报管理&#xff0c;新闻信息管理&#xff0c;维修人员管理 微信端账号功能包括&#xff1a;系统首页&#xff0c;新闻信息…
阅读更多...
C++ 视觉开发 六.特征值匹配

C++ 视觉开发 六.特征值匹配

以图片识别匹配的案例来分析特征值检测与匹配方法。 目录 一.感知哈希算法(Perceptual Hash Algorithm) 二.特征值检测步骤 1.减小尺寸 2.简化色彩 3.计算像素点均值 4.构造感知哈希位信息 5.构造一维感知哈希值 三.实现程序 1.感知哈希值计算函数 2.计算距离函数 3…
阅读更多...
CTF入门知识点

CTF入门知识点

CTF知识点 md5函数 <?php$a 123;echo md5($a,true); ?> 括号中true显示输出二进制 替换成false显示输出十六进制绕过 ffifdyop 这个字符串被 md5 哈希了之后会变成 276f722736c95d99e921722cf9ed621c&#xff0c;这个字符串前几位刚好是 or 6 而 Mysql 刚好又会把 …
阅读更多...
分支与循环

分支与循环

目录 1. if语句 1&#xff09;if 2) else 3&#xff09;分支中包含多条语句 4&#xff09;if嵌套 2.关系操作符 3.条件操作符 4.逻辑操作符&#xff1a;&& || ! 1) 逻辑取反运算符 !​编辑 2 与运算符​编辑 3) 或运算符​编辑 4) 闰年的判断 5) 短路 …
阅读更多...
最新文章
推荐文章

Copyright @ 2022~2023