CUDA技术测试-归约算法

最新推荐文章于 2025-03-10 16:40:30 发布
最新推荐文章于 2025-03-10 16:40:30 发布
阅读量271 收藏 1
点赞数
分类专栏： 并行计算（CUDA/HIP/多机多卡） 文章标签： cuda 并行计算
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.youkuaiyun.com/qq_33562636/article/details/109044798
版权 
#include "../common/common.h"
#include <cuda_runtime.h>
#include <stdio.h>

/*
 outdata=indata   指针复制,此语句使oudata指向的地址和indata指向的地址相同，它们都指向同一内存区域
 *outdata=*indata 指针赋值,此语句使指针outdata指向的地址（内存位置）的内容与指针indata指向地址（内存位置）的内容相同，但outdata与indata指向的地址不一定相
*/

// .execfilename ikernel arraysize -v (verify)  -p(print)
// ./matrix 1 32 1 1024 -v -p


void initialArray(float *array, int size)
{
	for (int i = 0; i < size; i++)
	{
		array[i] = (float)(rand() % 10 + 1);
	}
}


void printArray(const char *name, float *array, int size)
{
	printf("%s:arraySize:%d\n", name, size);
	for (int i = 0; i < size; i++)
	{
		printf("%4.0f", array[i]);
	}
	printf("\n");
}


int addblockSum(float *array, int size)
{
	float sum = 0.0;
	for (int i = 0; i < size; i++)
	{
		sum += array[i];
	}
	return sum;
}


int verifyArray(float sum1, float sum2)
{
	if (sum1 != sum2)
	{
		printf("\n The result %lf and %lf is not same!\n", sum1, sum2);
		return 1;
	}
	printf("\nThe result is same!\n");
	return 0;
}


/*
int recursiveReduce(float *indata,int const size)
{
	if (size == 1)
	{
	return indata[0];
	}

	int const stride = size / 2;

	for (int i = 0; i < stride; i++)
	{
		indata[i] += indata[i + stride];
	}
	return recursiveReduce(indata,stride);
}
*/


int recursiveReduce(float *indata, int const size)
{
	float sum = 0.0;
	for (int i = 0; i < size; i++)
	{
		sum += indata[i];
	}
	return sum;
}


//kernel 1

__global__ void reduceNeighbored(float *g_idata, float *g_odata, unsigned int size)
{
	unsigned int tid = threadIdx.x;
	unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;

	float *idata = g_idata + blockIdx.x * blockDim.x;

	if (idx >= size) return;

	for (int stride = 1; stride < blockDim.x; stride *= 2)
	{
		if ((tid % (2 * stride)) == 0)
		{
			idata[tid] += idata[tid + stride];
		}
		__syncthreads();
	}
	if (tid == 0) g_odata[blockIdx.x] = idata[0];
}


//kernel 2  邻近配对 block内的线程数量减少一半

__global__ void reduceNeighboredLess(float *g_idata, float *g_odata, unsigned int n)
{
	unsigned int tid = threadIdx.x;
	unsigned int idx = blockIdx.x * blockDim.x * 2 + threadIdx.x;

	float *idata = g_idata + blockIdx.x * blockDim.x * 2;

	if (idx >= n) return;

	for (int stride = 1; stride < blockDim.x * 2; stride *= 2)
	{
		int index = 2 * stride * tid;
		if (index < blockDim.x * 2)
		{
			idata[index] = idata[index] + idata[index + stride];
		}
		__syncthreads();
	}

	if (tid == 0) g_odata[blockIdx.x] = idata[0];
}


//kernel 3 交错配对 block内的线程数量减少一半

__global__ void reduceInterleaved(float *g_idata, float *g_odata, unsigned int n)
{
	unsigned int tid = threadIdx.x;
	unsigned int idx = blockIdx.x * blockDim.x * 2 + threadIdx.x;

	float *idata = g_idata + blockIdx.x * blockDim.x * 2;

	if (idx >= n) return;

	for (int stride = blockDim.x * 2 / 2; stride > 0; stride >>= 1)
	{
		if (tid < stride)
		{
			idata[tid] += idata[tid + stride];
		}
		__syncthreads();
	}

	if (tid == 0) g_odata[blockIdx.x] = idata[0];
}


//kernel 4  block数量减少一半

__global__ void reduceUnrolling2(float *g_idata, float *g_odata, unsigned int n)
{
	unsigned int tid = threadIdx.x;
	unsigned int idx = blockIdx.x * blockDim.x * 2 + threadIdx.x;

	float *idata = g_idata + blockIdx.x * blockDim.x * 2;

	if (idx + blockDim.x < n) g_idata[idx] += g_idata[idx + blockDim.x];

	__syncthreads();

	for (int stride = blockDim.x / 2; stride > 0; stride >>= 1)
	{
		if (tid < stride)
		{
			idata[tid] += idata[tid + stride];
		}
		__syncthreads();
	}
	if (tid == 0) g_odata[blockIdx.x] = idata[0];
}


//kernel 5  block数量减少到四分之一

__global__ void reduceUnrolling4(float *g_idata, float *g_odata, unsigned int n)
{
	unsigned int tid = threadIdx.x;
	unsigned int idx = blockIdx.x * blockDim.x * 4 + threadIdx.x;

	float *idata = g_idata + blockIdx.x * blockDim.x * 4;

	if (idx + 3 * 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];
		g_idata[idx] = a1 + a2 + a3 + a4;
	}

	__syncthreads();

	for (int stride = blockDim.x / 2; stride > 0; stride >>= 1)
	{
		if (tid < stride)
		{
			idata[tid] += idata[tid + stride];
		}
		__syncthreads();
	}
	if (tid == 0) g_odata[blockIdx.x] = idata[0];
}


//kernel 6  block数量减少到八分之一

__global__ void reduceUnrolling8(float *g_idata, float *g_odata, unsigned int n)
{
	unsigned int tid = threadIdx.x;
	unsigned int idx = blockIdx.x * blockDim.x * 8 + threadIdx.x;

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

	for (int stride = blockDim.x / 2; stride > 0; stride >>= 1)
	{
		if (tid < stride)
		{
			idata[tid] += idata[tid + stride];
		}
		__syncthreads();
	}
	if (tid == 0) g_odata[blockIdx.x] = idata[0];
}



// kernel 7  block数量减少到八分之一 线程束展开

__global__ void reduceUnrollWarps8(float *g_idata, float *g_odata, unsigned int n)
{
	unsigned int tid = threadIdx.x;
	unsigned int idx = blockIdx.x * blockDim.x * 8 + threadIdx.x;

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

	for (int stride = blockDim.x / 2; stride > 32; stride >>= 1)
	{
		if (tid < stride)
		{
			idata[tid] += idata[tid + stride];
		}
		__syncthreads();
	}

	if (tid < 32)
	{
		volatile float *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];
}


// kernel 8  block数量减少到八分之一 block内循环展开 完全循环展开

__global__ void reduceCompleteUnrollWarps8(float *g_idata, float *g_odata, unsigned int n)
{
	unsigned int tid = threadIdx.x;
	unsigned int idx = blockIdx.x * blockDim.x * 8 + threadIdx.x;

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




// kernel 9  block数量减少到八分之一 block内循环展开 完全循环展开 共享内存

__global__ void reduceSmemUnrolling4(float *g_idata, float *g_odata, unsigned int n)
{
	__shared__ float smem[1024];
	unsigned int tid = threadIdx.x;
	unsigned int idx = blockIdx.x * blockDim.x * 4 + threadIdx.x;

	float tmpSum = 0.0;

	if (idx + 3 * 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];
		tmpSum = a1 + a2 + a3 + a4;
	}

	smem[tid] = tmpSum;
	__syncthreads();

	if (blockDim.x >= 1024 && tid < 512) smem[tid] += smem[tid + 512];

	__syncthreads();

	if (blockDim.x >= 512 && tid < 256) smem[tid] += smem[tid + 256];

	__syncthreads();

	if (blockDim.x >= 256 && tid < 128) smem[tid] += smem[tid + 128];

	__syncthreads();

	if (blockDim.x >= 128 && tid < 64) smem[tid] += smem[tid + 64];

	__syncthreads();

	if (tid < 32)
	{
		volatile float *vsmem = smem;
		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] = smem[0];
}



int main(int argc, char **argv)
{
	int dev = 0;
	cudaDeviceProp deviceProp;
	cudaGetDeviceProperties(&deviceProp, dev);
	printf("%s starting transpose at ", argv[0]);
	printf("device %d: %s \n", dev, deviceProp.name);
	cudaSetDevice(dev);

	int arraySize = 1024;
	int dimx = 1024;
	int dimy = 1;

	int  ikernel = 0;
	bool ifverify = 0;
	bool ifprintf = 0;

	if (argc > 1) ikernel = atoi(argv[1]);
	if (argc > 2) dimx = atoi(argv[2]);
	if (argc > 3) dimy = atoi(argv[3]);
	if (argc > 4) arraySize = atoi(argv[4]);
	if ((argc > 5 && strcmp(argv[5], "-v") == 0))  ifverify = 1;
	if ((argc > 6 && strcmp(argv[6], "-p") == 0) || (argc > 5 && strcmp(argv[5], "-p") == 0))  ifprintf = 1;

	dim3 block(dimx, dimy);
	dim3 block2(dimx / 2, dimy);
	dim3 grid((arraySize + block.x - 1) / block.x, 1);
	dim3 grid4(grid.x / 2, 1);
	dim3 grid5(grid.x / 4, 1);
	dim3 grid6(grid.x / 8, 1);

	float *h_idata, *d_result;
	h_idata = (float*)malloc(arraySize * sizeof(float));
	d_result = (float*)malloc(grid.x * sizeof(float));

	initialArray(h_idata, arraySize);
	if (ifprintf) printArray("h_idata", h_idata, arraySize);

	float *d_idata, *d_odata;
	cudaMalloc((void**)&d_idata, arraySize * sizeof(float));
	cudaMalloc((void**)&d_odata, grid.x * sizeof(float));
	cudaMemcpy(d_idata, h_idata, arraySize * sizeof(float), cudaMemcpyHostToDevice);

	double iStart, iElaps;
	iStart = seconds();
	float cpu_sum = recursiveReduce(h_idata, arraySize);
	iElaps = seconds() - iStart;
	printf("cpu_sum:\n%0.0f\n", cpu_sum);
	printf("CPU运行时间为：%lfs\n\n", iElaps);

	char const *kernelName;
	iStart = seconds();
	switch (ikernel)
	{
	case 1:
		kernelName = "multiplicateMatrixOnDevice";
		reduceNeighbored << <grid, block >> > (d_idata, d_odata, arraySize);
		break;
	case 2:
		kernelName = "reduceNeighboredLess";
		block = block2;
		reduceNeighboredLess << <grid, block >> > (d_idata, d_odata, arraySize);
		break;
	case 3:
		kernelName = "reduceInterleaved";
		block = block2;
		reduceInterleaved << <grid, block >> > (d_idata, d_odata, arraySize);
		break;
	case 4:
		kernelName = "reduceUnrolling2 ";
		grid = grid4;
		reduceUnrolling2 << <grid, block >> > (d_idata, d_odata, arraySize);
		break;
	case 5:
		kernelName = "reduceUnrolling4 ";
		grid = grid5;
		reduceUnrolling4 << <grid, block >> > (d_idata, d_odata, arraySize);
		break;
	case 6:
		kernelName = "reduceUnrolling8 ";
		grid = grid6;
		reduceUnrolling8 << <grid, block >> > (d_idata, d_odata, arraySize);
		break;
	case 7:
		kernelName = "reduceUnrollWarps8 ";
		grid = grid6;
		reduceUnrollWarps8 << <grid, block >> > (d_idata, d_odata, arraySize);
		break;
	case 8:
		kernelName = "reduceCompleteUnrollWarps8 ";
		grid = grid6;
		reduceCompleteUnrollWarps8 << <grid, block >> > (d_idata, d_odata, arraySize);
		break;
	case 9:
		kernelName = "reduceSmemUnrolling4 ";
		grid = grid5;
		reduceSmemUnrolling4 << <grid, block >> > (d_idata, d_odata, arraySize);
		break;
	}

	iElaps = seconds() - iStart;
	cudaDeviceSynchronize();
	cudaMemcpy(d_result, d_odata, grid.x * sizeof(float), cudaMemcpyDeviceToHost);
	float gpu_sum = addblockSum(d_result, grid.x);

	printf("%s functions <<<grid(%d,%d),block(%d,%d)>>>  GPU运行时间为：%fs\n", kernelName, grid.x, grid.y, block.x, block.y, iElaps);
	printf("gpu_sum:\n%0.0f\n", gpu_sum);
	if (ifverify)	verifyArray(cpu_sum, gpu_sum);

	cudaFree(d_idata);
	cudaFree(d_odata);

	free(h_idata);
	free(d_result);

	return 0;
}

在这里插入图片描述

cuda编程---二维图像实现并行归约算法及基础优化方式
puqian13的博客
04-30 292
cuda编程---二维图像实现并行归约算法及基础优化方式 注： 1、本文不做任何公式推导，主要包含对算法的个人理解及少量的代码 2、以求取最大值为例，给出二维图像并行规约代码 代码中对于图像数据的处理，采用1维的共享内存和2维的共享内存分别实现归约 3、对于归约算法基础的优化方式（这里指连续寻址、分支发散），分析其背后的机制
CUDA 归约求和算法（一）
ikun的博仔的博客
05-16 3801
CUDA 归约求和算法： 并行归约（Reduction）是一种很基础的并行算法，简单来说，我们有N个输入数据，使用一个符合结合律的二元操作符作用其上，最终生成1个结果。这个二元操作符可以是求和、取最大、取最小、平方、逻辑与或等等。 由于加法的交换律和结合律，数组可以以任意顺序求和。所以我们会自然而然产生这样的思路：首先把输入数组划分为更小的数据块，之后用一个线程计算一个数据块的部分和，最后把所有部分和再求和得出最终结果。 上面演示了单个线程块内的归约求和过程，结合下面的代码解释： 13行第一轮循环，对应
基于CUDA的数组规约优化
12-17
细节参考：http://blog.youkuaiyun.com/qq_17239003/article/details/78824004#0-qzone-1-33634-d020d2d2a4e8d1a374a433f596ad1440
CUDA学习（归约算法）
tian_ciomp的博客
12-28 5887
parallel reduction Nvidia官网归约ppt 中文ppt 可以理解为将数组中所有数求和的过程并行化 CUDA本身并不支持全局同步，将每一层归约作为一个kernel重复递归调用
CUDA编程(2): 线程束执行的本质、规约问题、循环展开
最新发布
cumtchw
03-10 1013
CUDA编程(2): 线程束执行的本质、规约问题、循环展开
Cuda归约运算（一）——分析cudasample的reduction1
qq_32417133的博客
03-03 943
何谓归约 归约（redution）是一类并行算法，对传入的O(N)个输入数据，使用一个二元的复合结合律的操作符，生成O(1)的结果。这类操作包括取最小、最大、求和、平方求和、逻辑与、逻辑或、向量点积。归约也是其他高级运算中要用的基础算法。 除非操作符的求解代价极高，否则归约倾向于带宽受限型任务。下面就从SDK提供的reduction例子入手，详细理解该归约算法。 概述 因为该二元操作符复合结合律...
CUDA并行归约算法（一）
RandyOliver's House
07-11 1084
并行归约（Reduction）是一种并行算法，对于符合结合律的二元操作符，将输入的数组划分为更小的数据库，每个线程计算1个数据块的部分结果，最后把所有部分结果再计算，得出最终结果。二元操作符可以是求和、取最大、取最小、平方、逻辑与或等。
三种基于CUDA的归约计算
yan31415的博客
03-02 2245
归约在并行计算中很常见，并且在实现上具有一定的套路。本文分别基于三种机制（Intrinsic，共享内存，atomic），实现三个版本的归约操作，完成一个warp大小的数据的归约求和计算。 Intrinsic版本 基于Intrinsic函数 __shfl_down_sync 实现，使一个warp内的线程通过读取相邻线程寄存器中数据，完成归约操作。 实现如下： __global__ void kIntrinsicWarpReduce(int* d_src, int* d_dst) { int val = d_
cuda并行归约算法（不断优化过程）
weixin_45806592的博客
09-13 1304
归约
CUDA编程练习（九） 归约
SimpleBright_Man的博客
06-03 579
拖延症晚期，但是还有得救
CUDA----规约
UCAS_HMM的博客
10-18 1185
CUDA规约操作
CUDA的归约及求和优化
Begoniaish的博客
04-14 1067
（3） 对每种实现，重复实验2000次，统计平均时间用于最终性能评价（对于任意长度的多块扫描，尝试不同的线程块大小）。8种实现方式，完成PPT里所提及的思考题；（2） 利用thrust库实现并行扫描；消除bank conflict。单线程加载全局内存时做一次加法。单线程加载全局内存时做多次加法。消除bank conflict。单线程加载全局内存时做一次加法。单线程加载全局内存时做多次加法。shuffle指令优化。shuffle指令优化。
CUDA并行算法系列之规约
張行之的专栏
09-21 6064
CUDA并行算法系列之规约前言规约是一类并行算法，对传入的N个数据，使用一个二元的符合结合律的操作符⊕，生成1个结果。这类操作包括取最小、取最大、求和、平方和、逻辑与/或、向量点积。规约也是其他高级算法中重要的基础算法。除非操作符⊕的求解代价极高，否则规约倾向于带宽受限型任务（bandwidthbound）。本文将介绍几种规约算法的实现，从两遍规约、block的线程数必须为2的幂，一步一步优化。
CUDA处理归约问题
yan31415的博客
11-20 535
归约问题由于计算操作的对称性，非常适合并行处理。本文以数组求和为例，通过CUDA先实基础版本，并基于基础版本尝试通过不同的优化手段实现几个方案，最后将所有优化手段集成到最终的实现。 其中核函数执行时间是通过Nsight System工具抓取，命令参数如下 nsys profile --stats true a.out 5 基础版本实现 核函数实现 __global__ void kReduceWholeWarpSafe(int64_t* d_src, int64_t* d_sum) { int64_t*
CUDA高性能计算经典问题①：归约
OneFlow深度学习框架
10-29 1277
撰文 |Will Zhang本系列为CUDA进阶，通过具体的经典问题，讲述高性能编程的一些基本原则以及方法。建议读者先阅读NVIDIA官方的编程指南完成CUDA入门，基础比较少的同学也建...
CUDA中并行规约（Parallel Reduction）的优化
weixin_30374009的博客
11-08 610
转自：http://hackecho.com/2013/04/cuda-parallel-reduction/ Parallel Reduction是NVIDIA-CUDA自带的例子，也几乎是所有CUDA学习者的的必看算法。在这个算法的优化中，Mark Harris为我们实现了7种不同的优化版本，将Bandwidth几乎提高到了峰值。相信我们通过仔细研读这个过程，一定能对CUDA程序的优化有更...
CUDA编程学习＜2＞——归约算法的7种优化方法详解
静以修身，简以养德
02-16 2374
参考链接： CUDA学习（归约算法） 根据上面文章中的PPT链接，进行了深入研究。原文章只介绍了归约算法的一般形式，而PPT中则对比了几种归约算法，但是PPT写的不够完整，不便于系统的运行起来。为了学习CUDA及其优化技巧，特此整理 Nvidia官网归约ppt （英文） 中文PPT 关于代码核函数的部分可以参阅我的博客： CUDA编程学习——关于核函数ID获取公式 硬件信息 实验机器硬件： 通过硬件信息可以看出几个重要信息： 理论带宽(bandwidth)：192GB/s 此显卡每秒渲染像素数量（Pixe
CUDA中的归约
weixin_34186931的博客
12-29 342
CUDA编程实战书中的乘方和解决办法： 对一个数组执行某种计算，然后产生一个更小的结果数组。 由一个线程在共享内存上进行迭代并计算出总和值。而如果用并行，所花时间就与数组长度的对数成正比。 代码的思想是，每个线程将cache[]中的两个值加和，然后结果保存回cache[]。由于每个线程都将两个值合并为一个值，那么在完成这个步骤后，得到的结果数量就是计算开始时数值数量的一半。在下一个步骤中，我...
CUDA并行计算课程作业 - CUDA Scan算法详解
实现CUDA Scan算法时，需要考虑几个关键的技术点： 1. 内存访问模式：由于GPU的架构，内存访问模式对性能有巨大影响。好的内存访问模式可以充分利用内存带宽，而坏的内存访问模式会因为内存访问冲突导致性能下降。 ...
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值