CUDA技术测试-归约算法

最新推荐文章于 2023-12-07 15:22:12 发布
最新推荐文章于 2023-12-07 15:22:12 发布
阅读量229 收藏 1
点赞数
分类专栏： 并行计算（CUDA/HIP/多机多卡） 文章标签： cuda 并行计算
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/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;
}

在这里插入图片描述

Rocky_myx
关注
  • 0
    点赞
  • 踩
  • 1
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
专栏目录
CUDA编程学习＜2＞——归约算法的7种优化方法详解
静以修身，简以养德
02-16 1417
参考链接： CUDA学习（归约算法） 根据上面文章中的PPT链接，进行了深入研究。原文章只介绍了归约算法的一般形式，而PPT中则对比了几种归约算法，但是PPT写的不够完整，不便于系统的运行起来。为了学习CUDA及其优化技巧，特此整理 Nvidia官网归约ppt （英文） 中文PPT 关于代码核函数的部分可以参阅我的博客： CUDA编程学习——关于核函数ID获取公式 硬件信息 实验机器硬件： 通过硬件信息可以看出几个重要信息： 理论带宽(bandwidth)：192GB/s 此显卡每秒渲染像素数量（Pixe
CUDA程序性能优化 并行归约
weixin_62298882的博客
03-21 122
基本思想是，对于一个输入数组执行某种计算，然后产生一个更小的结果数组。当大量的数进行加和运算时，可以利用归约算法，多线程进行求和运算。解决方法: 相邻线程执行相同的代码。成对方式导致线程发散如下。​​​​​​归约算法。
CUDA----规约
UCAS_HMM的博客
10-18 759
CUDA规约操作
CUDA规约求和/最大值
保持分享欲
12-07 1608
即得到最后结果。如下图所示：值得注意的是，最后的多个规约结果可以拷贝到CPU中计算，也可以在GPU进行计算。在GPU中进行计算时，当规约结果数量小于单个线程块的线程数时，即可得到最终结果。
GPU编程 CUDA C++ 数组归约的示例和解释
行子的博客
03-31 376
在主函数中，我们首先定义了输入数组的大小（即“size”）、每个线程块中的线程数（即“threads_per_block”）和线程块的数量（即“blocks”）。然后，如果当前线程ID小于“s”，则将当前线程的共享内存值与其相邻的线程的共享内存值相加。在此示例中，我们使用了一些简单的技巧来优化内核函数的性能，例如使用共享内存来提高内存访问效率，并使用循环执行归约操作。最后，我们使用“cudaMemcpy”函数将输出数组从设备内存复制到主机内存，并计算输出数组中的所有元素的和。
CUDA 并行规约问题的普通实现以及嵌套递归实现（附加嵌套递归的一点小问题解决方法）
Coco9_night的博客
06-03 1801
规约问题串行规约问题问题分析串行代码实现过程总代码并行规约关于GPU的grid和block的概念普通并行规约问题普通并行代码简单实现过程普通并行代码的无耦合(关于grid和block的耦合)实现解决一些小问题嵌套递归并行规约问题GPU动态并行动态并行在规约中的使用完整代码（包括main函数）动态并行出现的问题感谢阅读 串行规约问题 问题分析 首先来分析一个简单的规约问题，假设我们此时有一个简单数组 size=8; int array[size]={1, 2, 3, 4, 5, 6, 7, 8}; 我们的目
CUDA_GPU_前缀求和算法.cu
05-18
基于CUDA的GPU并行，给出了三种不同的前缀求和算法，第一种是基本的规约并行算法，第二种是采用共享内存优化的GPU并行算法，第三种是采用trust库的前缀求和函数。并且给出了三种方法对比之间的性能差异
cuda-使用cuda并行加速实现之reduction.zip
最新发布
04-24
CUDA（Compute Unified Device Architecture）是由NVIDIA推出的编程模型，它允许...本资料中的“cuda_使用cuda并行加速实现之reduction”应涵盖了这些概念，并可能提供了实际的代码示例，帮助理解并实现CUDA归约。
面向GPU计算平台的归约算法的性能优化研究.pdf
09-25
《面向GPU计算平台的归约算法的性能优化研究》是一篇探讨如何提高GPU计算平台上的归约算法效率的学术论文。归约算法在科学计算、图像处理等多个领域有着广泛的应用，是并行计算的核心算法之一。因此，对其进行优化以...
CUDA.rar_cuda process_threads
09-14
在实现CUDA归约过程中，有几种经典的算法策略，如ShuffleReduce、WarpReduce和BlockReduce等。这些算法考虑了线程之间的通信和同步，以及如何有效地利用共享内存。例如，WarpReduce是利用CUDA硬件中的 warp-level ...
Optimizing parallel reduction in CUDA 规约优化文档
12-25
Optimizing parallel reduction in CUDA 规约优化文档
CUDA专家手册 GPU编程权威指南
02-24
《CUDA专家手册：GPU编程权威指南》详细讨论CUDA的硬件...　第三部分是案例剖析，深入分析精选的CUDA应用场景以及关键的并行算法，包括流式负载、归约、扫描（并行前缀求和）、N-体问题和图像处理，这些算法全方位涵盖
CUDA并行算法系列之规约
張行之的专栏
09-21 5693
CUDA并行算法系列之规约前言规约是一类并行算法，对传入的N个数据，使用一个二元的符合结合律的操作符⊕，生成1个结果。这类操作包括取最小、取最大、求和、平方和、逻辑与/或、向量点积。规约也是其他高级算法中重要的基础算法。除非操作符⊕的求解代价极高，否则规约倾向于带宽受限型任务（bandwidthbound）。本文将介绍几种规约算法的实现，从两遍规约、block的线程数必须为2的幂，一步一步优化。
CUDA并行规约(交错配对-完全展开-终极版）
武泗海的博客
12-17 1229
通过前面的文章，我们对32以下的迭代循环进行了展开处理，实际上，由于线程块的长度限制（GTX1050Ti是1024），可以说循环次数是确定的，因此可以将循环完全展开， 即 1024 512 256 128 64 都可以展开计算，唯一需要注意的每次计算之后都要进行同步，原因在前文已有解释。下面给出代码。 #include "cuda_runtime.h" #include "device_lau
CUDA ---- Branch Divergence and Unrolling Loop
weixin_30583563的博客
06-02 281
Avoiding Branch Divergence 有时，控制流依赖于thread索引。同一个warp中，一个条件分支可能导致很差的性能。通过重新组织数据获取模式可以减少或避免warp divergence（该问题的解释请查看warp解析篇）。 The Parallel Reduction Problem 我们现在要计算一个数组N个元素的和。这个过程用CPU编程很容易实现： int ...
CUDA学习笔记（七）分支发散和展开循环
云中寻雾的博客
07-20 3713
这节是关于Branch Divergence and Unrolling Loop。（分支发散和展开循环） Avoiding Branch Divergence（避免分支发散） 有时，控制流依赖于thread索引。同一个warp中，一个条件分支可能导致很差的性能。通过重新组织数据获取模式可以减少或避免warp divergence（该问题的解释请查看warp解析篇）。 The Paralle...
cuda编程_CUDA编程入门（五）更高效的并行归约算法
weixin_39682301的博客
11-29 135
上一篇我们写了一个简单的并行归约算法，虽然相比与CPU已经达到了较快的速度，但其实程序还有很大的优化空间。比较明显的是有两个需要优化的点：线程束分化内存访问下面我们来一一解决。线程束分化首先我们来回忆一下上个程序的实现思路，我们是利用stride变量来实现每轮迭代时的被加数选择，这造成线程束的分化，也就是每个线程束中只有部分线程是活跃的，但由于硬件设计，调度会以一整个线程束为单位进行，所以影响了程...
CUDA系列学习（五）GPU基础算法: Reduce, Scan, Histogram
热门推荐
Rachel Zhang的专栏
06-25 4万+
喵~不知不觉到了CUDA系列学习第五讲，前几讲中我们主要介绍了基础GPU中的软硬件结构，内存管理，task类型等；这一讲中我们将介绍3个基础的GPU算法：reduce，scan，histogram，它们在并行算法中非常常用，我们在本文中分别就其功能用处，串行与并行实现进行阐述。 1. Task complexitytask complexity包括step complexity（可以并行成几个操作

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

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

抵扣说明：

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

余额充值