【CUDA】 扫描 Scan

news2024/11/17 2:54:43

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 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
  
    // Write result 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];
}

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 memory
    XY[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 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];
}

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 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();

    // Phase 2
    // Use Brent-Kung algorithm to scan the last elements of each section
    for (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 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];
}

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

    // Free memory on the device
    cudaFree(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.coloradmin.cn/o/1900816.html

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

相关文章

Android Graphics 显示系统 - 监测、计算FPS的工具及设计分析

“ 在Android图像显示相关的开发、调试、测试过程中&#xff0c;如何能有效地评估画面的流畅度及监测、计算图层渲染显示的实时FPS呢&#xff1f;本篇文章将会提供一种实用、灵巧的思路。” 01 设计初衷 面对开发测试中遇到的卡顿掉帧问题&#xff0c;如何在复现卡顿的过程中持…

黑马的ES课程中的不足

在我自己做项目使用ES的时候&#xff0c;发现了黑马没教的方法&#xff0c;以及一些它项目的小问题 搜索时的匹配方法 这个boolQuery().should 我的项目是通过文章的标题title和内容content来进行搜索 但是黑马它的项目只用了must 如果我们的title和content都用must&#x…

QCustomPlot+ vs2022+ qt

零、printSupport 步骤一&#xff1a;下载QCustomPlot 访问QCustomPlot的官网 QCustomPlot 下载最新版本的源代码。 步骤二&#xff1a;配置项目 创建新的Qt项目&#xff1a; 打开VS2022&#xff0c;创建一个新的Qt Widgets Application项目。 将QCustomPlot源代码添加到项目…

C语言编程与进阶

1.0 C语言关键字 1-1C语言关键字-CSDN博客文章浏览阅读831次&#xff0c;点赞13次&#xff0c;收藏24次。define使用define定义常量return 0;使用define定义宏// define 定义宏&#xff0c;名字是ADD(x,y),x y 是宏的参数int a 10;int b 20;return 0;宏定义的本质是替换&am…

JavaEE——计算机工作原理

冯诺依曼体系&#xff08;VonNeumannArchitecture&#xff09; 现代计算机&#xff0c;大多遵守冯诺依曼体系结构 CPU中央处理器&#xff1a;进行算术运算与逻辑判断 存储器&#xff1a;分为外存和内存&#xff0c;用于存储数据&#xff08;使用二进制存储&#xff09; 输入…

百日筑基第十二天-入门Elasticsearch

百日筑基第十二天-入门Elasticsearch Elasticsearch 是什么 Elasticsearch 是一个分布式、RESTful 风格的搜索和数据分析引擎。 安装 Elasticsearch 下载&#xff1a;https://www.elastic.co/cn/downloads/elasticsearch Elasticsearch 是免安装的&#xff0c;只需要把 zip…

绝了,华为伸缩摄像头如何突破影像边界?

自华为Pura70 Ultra超聚光伸缩镜头诞生以来&#xff0c;备受大家的关注&#xff0c;听说这颗镜头打破了传统手机的摄像头体积与镜头的设计&#xff0c;为我们带来了不一样的拍照体验。 智能手机飞速发展的今天&#xff0c;影像功能已经成为我们衡量一款手机性能的重要指标。想…

【Qt5.12.9】程序无法显示照片问题(已解决)

问题记录&#xff1a;Qt5.12.9下无法显示照片 我的工程名为03_qpainter&#xff0c;照片cd.png存放在工程目录下的image文件夹中。 /03_qpainter/image/cd.png 因为这是正点原子Linux下Qt书籍中的例程&#xff0c;在通过学习其配套的例程中的项目&#xff0c;发现我的项目少…

Python的招聘数据分析与可视化管理系统-计算机毕业设计源码55218

摘要 随着互联网的迅速发展&#xff0c;招聘数据在规模和复杂性上呈现爆炸式增长&#xff0c;对数据的深入分析和有效可视化成为招聘决策和招聘管理的重要手段。本论文旨在构建一个基于Python的招聘数据分析与可视化管理系统。 该平台以主流招聘平台为数据源&#xff0c;利用Py…

昇思25天学习打卡营第1天|初识MindSpore

# 打卡 day1 目录 # 打卡 day1 初识MindSpore 昇思 MindSpore 是什么&#xff1f; 昇思 MindSpore 优势|特点 昇思 MindSpore 不足 官方生态学习地址 初识MindSpore 昇思 MindSpore 是什么&#xff1f; 昇思MindSpore 是全场景深度学习架构&#xff0c;为开发者提供了全…

Ubuntu固定虚拟机的ip地址

1、由于虚拟机网络是桥接&#xff0c;所以ip地址会不停地变化&#xff0c;接下来我们就讲述ip如何固定 2、如果apt安装时报错W: Target CNF (multiverse/cnf/Commands-all) is configured multiple times in /etc/apt/sources.list:10&#xff0c; 检查 /etc/apt/sources.list…

计算机组成原理--概述

&#x1f308;个人主页&#xff1a;小新_- &#x1f388;个人座右铭&#xff1a;“成功者不是从不失败的人&#xff0c;而是从不放弃的人&#xff01;”&#x1f388; &#x1f381;欢迎各位→点赞&#x1f44d; 收藏⭐️ 留言&#x1f4dd; &#x1f3c6;所属专栏&#xff1…

AI Earth应用—— 在线使用sentinel数据VV和VH波段进行水体提取分析(昆明抚仙湖、滇池为例)

AI Earth 本文的主要目的就是对水体进行提取,这里,具体的操作步骤很简单基本上是通过,首页的数据检索,选择需要研究的区域,然后选择工具箱种的水体提取分析即可,剩下的就交给阿里云去处理,结果如下: 这是我所选取的一景影像: 详情 卫星: Sentinel-1 级别: 1 …

Redis IO多路复用

0、前言 本文所有代码可见 > 【gitee code demo】 本文涉及的主题&#xff1a; 1、BIO、NIO的业务实践和缺陷 2、Redis IO多路复用&#xff1a;redis快的主要原因 3、epoll 架构 部分图片 via 【epoll 原理分析】 1、BIO单线程版 1.1 业务代码 client client代码相同…

Proxmox VE 8虚拟机直通USB磁盘

作者&#xff1a;田逸&#xff08;fromyz&#xff09; 今天有个兄弟发消息&#xff0c;咨询怎么让插在服务器上的U盾被Proxmox VE上的虚拟机识别。在很久很久以前&#xff0c;我尝试过在Proxmox VE 5以前的版本创建windows虚拟机&#xff0c;并把插在Proxmox VE宿主机上的银行U…

Vue3基础(二)

一、搭建工程(vite) ## 1.创建命令 npm create vuelatest## 2.具体配置 ## 配置项目名称 √ Project name: vue3_test ## 是否添加TypeScript支持 √ Add TypeScript? Yes ## 是否添加JSX支持 √ Add JSX Support? No ## 是否添加路由环境 √ Add Vue Router for Single P…

【matlab】智能优化算法——基准测试函数

智能优化算法的基准测试函数是用于评估和优化算法性能的一组标准问题。这些测试函数模拟了真实世界优化问题的不同方面&#xff0c;包括局部最小值、全局最优解、高维度、非线性、不连续等复杂性。以下是对智能优化算法基准测试函数的详细归纳&#xff1a; 测试函数的分类&…

使用nohup和CUDA_VISIBLE_DEVICES进行GPU训练的教程

文章目录 1. 在单个GPU上训练模型1.1 使用nohup命令运行Python脚本1.2 查看运行中的进程1.3 查看输出日志 2. 在多个GPU上训练模型2.1 启动第一个程序&#xff0c;指定使用第0号GPU2.2 启动第二个程序&#xff0c;指定使用第1号GPU2.3 查看运行中的进程2.4 查看输出日志 3. 总结…

【不容错过】可灵AI重磅更新:画质升级,运镜控制,首尾帧自定义,还有30万创作激励奖金!

还记得最近在各大平台肆虐的老照片变成视频吗&#xff0c;就是用快手的可灵AI做的&#xff0c;今天可灵又迎来了一次重大更新。 「电脑端上线了」 之前一直用其他工具生的图片还需要保存到手机上&#xff0c;再用可灵来生成视频&#xff0c;很多人都能感受到手机操作不太方便&…

【ARMv8/v9 GIC 系列 5.6 -- GIC 超优先级中断详细介绍】

请阅读【ARM GICv3/v4 实战学习 】 文章目录 Interrupt superpriority超优先级中断的特性和应用Physical interface interrupt signalsPhysical Group 1 Non-NMI for Current Security StatePhysical Group 1 for Other Security State, or a Group 0 Non-NMIPhysical Group 1 …