归约求和程序，每次运行结果都不相同

laughcry2002 2010-01-21 08:14:01
刚学习cuda编程，试着写一个数组归约求和的程序，结果每次运行结果都不相同。猜测应该是哪里同步出问题了，可调试了很久都没找到原因。

平常很少拿问题来论坛来提问，可这个问题困扰我近一周时间了，不知哪位肯指点一下？多谢了。


---------- mini_test.cpp ------------


#include <cmath>

#include <cutil_inline.h>



typedef float Real;



#define BLOCK_SIZE 8



__device__ void 

kernel_vec_subsum(unsigned int nz, Real* g_idata, Real* g_odata)

{

    unsigned int tid = threadIdx.x;

    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;



	__shared__ Real sdata[BLOCK_SIZE];



	sdata[tid] = (i < nz) ? g_idata[i] : 0;

	__syncthreads();



	for(unsigned int s = BLOCK_SIZE / 2; s > 0; s >>= 1)

	{

		if (tid < s)	sdata[tid] += sdata[tid + s];

		__syncthreads();

	}



	__syncthreads();

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

}



__device__ void 

kernel_vec_sum(unsigned int nz, Real* g_vec, Real* g_sum)

{

    unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

	unsigned int size = nz;

	while( size > 1 ) 

	{

		kernel_vec_subsum(size, g_vec, g_vec);

		size = (size + BLOCK_SIZE - 1) / BLOCK_SIZE;

		__syncthreads();

	}



	if(i == 0)		g_sum[0] = g_vec[0];

}



__global__ void 

global_vec_sum(unsigned int nz, Real* g_vec, Real* g_sum)

{

	kernel_vec_sum(nz, g_vec, g_sum);

}



Real 

vec_sum(unsigned int nz, Real* vec)

{

	Real* g_idata;

	size_t mem_size = nz * sizeof(Real);

	cutilSafeCall( cudaMalloc( (void**) &g_idata, mem_size));

	cutilSafeCall( cudaMemcpy( g_idata, vec, mem_size, cudaMemcpyHostToDevice) );



	Real* g_odata;

	cutilSafeCall( cudaMalloc( (void**) &g_odata, sizeof(Real)));



	global_vec_sum <<< (nz + BLOCK_SIZE - 1) / (BLOCK_SIZE), BLOCK_SIZE >>>

		(nz, g_idata, g_odata);



	Real sum;

	cutilSafeCall( cudaMemcpy( &sum, g_odata, sizeof(Real), cudaMemcpyDeviceToHost) );



	cutilSafeCall(cudaFree(g_idata));

	cutilSafeCall(cudaFree(g_odata));



	return sum;

}



void test_vec_sum()

{

	unsigned int nz = 80 * BLOCK_SIZE + 3; // 80 * 8 + 3 = 643



	Real* vec = new Real[nz];

	for(unsigned int i = 0; i < nz; ++i)

		vec[i] = 1.0;



	Real sum = vec_sum(nz, vec);

	

    printf("%5.2f ", sum);



    delete[] vec;

}



int main( int argc, char** argv ) 

{

	cutilDeviceInit(argc, argv);

	test_vec_sum();

	cutilExit(argc, argv);

}
...全文
829 13 打赏 收藏 转发到动态 举报
写回复
用AI写文章
13 条回复
切换为时间正序
请发表友善的回复…
发表回复
lhw978 2011-06-09
  • 打赏
  • 举报
 回复
块间同步，threadfence确实不是同步!。
tanqiang917 2010-08-09
  • 打赏
  • 举报
 回复
我只想看下答案，至于还得注册回复吗!!
laughcry2002 2010-01-23
  • 打赏
  • 举报
 回复
[Quote=引用 10 楼 l7331014 的回复:]
对超过一个block的规约问题,就该通过反复启动kernel来同步的.呵呵.
当然,为更好的性能,要尽量启动一次kernel来处理最大的"块".(同一block中可以通过__syncthreads();同步的).

另外,threadfence不是同步!

下面是樟树同志的解说:
threadfence不是保证所有线程都完成同一操作
而只保证正在进行fence的线程本身的操作能够对所有线程安全可见
fence不要求线程运行到同一指令，而barrier有要求
[/Quote]

原来如此，现在终于明白了。我读cuda的编程指南，感觉真是惜墨如金、字字珠玑、点到为止啊，例子都不肯举一些（比如threadfence，确实手册上是说明了其作用是“让该句之前对global/shared内在的访问操作能够被所有其他线程可见”，读时直觉就是一种同步，呵，结果就犯错误）。

多亏有像本论坛这样的课外讨论与交流的平台啊。

再次感谢17331014君的热心帮助！

帖子虽然结了，仍然欢迎大家多指教。
laughcry2002 2010-01-22
  • 打赏
  • 举报
 回复
顺便问一下，发贴时代码部分难道不是用[-code-] [-/code-] （去除减号）括起来吗？怎么上上楼代码部分不见了。呵
laughcry2002 2010-01-22
  • 打赏
  • 举报
 回复
仍然有问题：

我试着将上述的 __device__ 函数 kernel_vec_sum 修改如下（与之相适应，调用此函数时要为 g_sum 分配 nz 个 Real 的存储空间），结果仍然不确定。

__device__ void
kernel_vec_sum_v2(unsigned int nz, Real* g_vec, Real* g_sum)
// 这里要求 g_sum 是与 g_vec 等长的向量，用于临时交换数据
{
unsigned int tid = threadIdx.x;
unsigned int bid = blockIdx.x;
unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int size = nz;

while( size > 1 )
{
kernel_vec_subsum(size, g_vec, g_sum);
size = (size + BLOCK_SIZE - 1) / BLOCK_SIZE;
__syncthreads();

// 将本轮计算结果 从g_sum 复制到 g_vec 向量中
if( tid == 0 && bid < size )
g_vec[bid] = g_sum[bid];
__syncthreads();
}

}


而当将该函数中所用的流程用于 host 端的函数时（函数vec_sum_v2中直接循环调用global_vec_subsum ），则不会产生问题。仍然没有想到问题出在什么地方？


__global__ void
global_vec_subsum(unsigned int nz, Real* g_idata, Real* g_odata)
{
kernel_vec_subsum(nz, g_idata, g_odata);
}

Real
vec_sum_v2(unsigned int nz, Real* vec)
{
Real* g_idata;
size_t mem_size = nz * sizeof(Real);
cutilSafeCall( cudaMalloc( (void**) &g_idata, mem_size));
cutilSafeCall( cudaMemcpy( g_idata, vec, mem_size, cudaMemcpyHostToDevice) );

Real* g_odata;
cutilSafeCall( cudaMalloc( (void**) &g_odata, mem_size));

while( nz > 1 ) {
unsigned int nz_new = (nz + BLOCK_SIZE - 1) / BLOCK_SIZE;
global_vec_subsum <<< nz_new, BLOCK_SIZE >>>
(nz, g_idata, g_odata);

nz = nz_new;
cutilSafeCall( cudaMemcpy( g_idata, g_odata, nz * sizeof(Real), cudaMemcpyDeviceToDevice) );
}

Real sum;
cutilSafeCall( cudaMemcpy( &sum, g_odata, sizeof(Real), cudaMemcpyDeviceToHost) );

cutilSafeCall(cudaFree(g_idata));
cutilSafeCall(cudaFree(g_odata));

return sum;
}

laughcry2002 2010-01-22
  • 打赏
  • 举报
 回复
仍然有问题：

我试着将上述的 __device__ 函数 kernel_vec_sum 修改如下（与之相适应，调用此函数时要为 g_sum 分配 nz 个 Real 的存储空间），结果仍然不确定。



__device__ void 

kernel_vec_sum_v2(unsigned int nz, Real* g_vec, Real* g_sum)

// 这里要求 g_sum 是与 g_vec 等长的向量，用于临时交换数据

{

	unsigned int tid = threadIdx.x;

	unsigned int bid = blockIdx.x;

	unsigned int i = blockIdx.x * blockDim.x + threadIdx.x;

	unsigned int size = nz;

	

	while( size > 1 ) 

	{

		kernel_vec_subsum(size, g_vec, g_sum);

		size = (size + BLOCK_SIZE - 1) / BLOCK_SIZE;

		__syncthreads();

		

		// 将本轮计算结果 从g_sum 复制到 g_vec 向量中

		if( tid == 0 && bid < size )

			g_vec[bid] = g_sum[bid];

		__syncthreads();

	}



}


而当将该函数中所用的流程用于 host 端的函数时（函数vec_sum_v2中直接循环调用global_vec_subsum ），则不会产生问题。仍然没有想到问题出在什么地方？


__global__ void 

global_vec_subsum(unsigned int nz, Real* g_idata, Real* g_odata)

{

	kernel_vec_subsum(nz, g_idata, g_odata);

}



Real 

vec_sum_v2(unsigned int nz, Real* vec)

{

	Real* g_idata;

	size_t mem_size = nz * sizeof(Real);

	cutilSafeCall( cudaMalloc( (void**) &g_idata, mem_size));

	cutilSafeCall( cudaMemcpy( g_idata, vec, mem_size, cudaMemcpyHostToDevice) );



	Real* g_odata;

	cutilSafeCall( cudaMalloc( (void**) &g_odata, mem_size));



	while( nz > 1 ) {

		unsigned int nz_new = (nz + BLOCK_SIZE - 1) / BLOCK_SIZE;

		global_vec_subsum <<< nz_new, BLOCK_SIZE >>>

			(nz, g_idata, g_odata);

		

		nz = nz_new;

		cutilSafeCall( cudaMemcpy( g_idata, g_odata, nz * sizeof(Real), cudaMemcpyDeviceToDevice) );

	}

	

	Real sum;

	cutilSafeCall( cudaMemcpy( &sum, g_odata, sizeof(Real), cudaMemcpyDeviceToHost) );



	cutilSafeCall(cudaFree(g_idata));

	cutilSafeCall(cudaFree(g_odata));



	return sum;

}


无心人_过过小日子 2010-01-22
  • 打赏
  • 举报
 回复
[Quote=引用 2 楼 laughcry2002 的回复:]
当线程号 i > nz 时，sdata[tid]取为0值；又因为是求和，取0值累加对结果没影响
[/Quote]

没多想,好像是如此.呵呵.

[Quote=引用 2 楼 laughcry2002 的回复:]
改用两个（一个用作 input, 另一个用作 output)
[/Quote]

应该可以解决问题吧.(性能不算)
无心人_过过小日子 2010-01-22
  • 打赏
  • 举报
 回复
对超过一个block的规约问题,就该通过反复启动kernel来同步的.呵呵.
当然,为更好的性能,要尽量启动一次kernel来处理最大的"块".(同一block中可以通过__syncthreads();同步的).

另外,threadfence不是同步!

下面是樟树同志的解说:
threadfence不是保证所有线程都完成同一操作
而只保证正在进行fence的线程本身的操作能够对所有线程安全可见
fence不要求线程运行到同一指令，而barrier有要求
laughcry2002 2010-01-22
  • 打赏
  • 举报
 回复
[Quote=引用 7 楼 l7331014 的回复:]
__syncthreads();
只同步一个block之内的,block之间的不同步!
因此上面的修改1等于无效.呵呵.
[/Quote]

这个我后来也注意到了，将 __syncthreads() 换为 __threadfence() 应该使得所有线程同步吗？但尝试的结果似乎仍然不行。

都快被这个程序折磨得没脾气了。我本意是想将 vec_sum(...) 包装成一个 __device__ 函数以便其他的 __device__ 或 __global__ 函数来调用它。举个例子，求两个向量内积的 __device__ vec_dot(...) 函数就可会用到这个向量元素求和的 vec_sum。

l7331014君还有什么建议提示一下？谢谢！
无心人_过过小日子 2010-01-22
  • 打赏
  • 举报
 回复
下一修改是正解.
无心人_过过小日子 2010-01-22
  • 打赏
  • 举报
 回复
__syncthreads();
只同步一个block之内的,block之间的不同步!
因此上面的修改1等于无效.呵呵.
laughcry2002 2010-01-21
  • 打赏
  • 举报
 回复
非常感谢l7331014的指导！

第1个问题我觉得不会有什么问题，因为 for 循环内操作的是 shared 内存变量 sdata，这个变量在参与计算的任何线程块中都是有值的（当线程号 i > nz 时，sdata[tid]取为0值；又因为是求和，取0值累加对结果没影响），从而不会产生问题。

您提到的第2个问题可能是出错的症结所在，由于手边没有调试环境，我明天上班时检验一下效果。初步设想解决的办法是调用 kernel_vec_subsum(size, g_vec, g_vec) 时后两个参数不能使用同一数组，改用两个（一个用作 input, 另一个用作 output）。不知您觉得如何？
无心人_过过小日子 2010-01-21
  • 打赏
  • 举报
 回复
好像有2个问题:
1)
for(unsigned int s = BLOCK_SIZE / 2; s > 0; s >>= 1)
{
if (tid < s) sdata[tid] += sdata[tid + s];
__syncthreads();
}
当nz为奇数时,结果正确吗?LZ再仔细考虑一下.
2)
kernel_vec_subsum(size, g_vec, g_vec);
假设g_vec被分成k份BLOCK_SIZE,"宏观"上这k份是并行计算,得到g_vec[h](h=0...k-1).而且在并行时,在"微观"上实际计算顺序又是不确定的,有可能在计算某一个份时,输入的已经其它的归约结果了....
go分布式mapreduce与raft算法与分布式数据库
掌握分布式mapreduce与raft算法与分布式数据库MapReduce是一种编程模型，用于大规模数据集（大于1TB）的并行运算。概念Map（映射）和Reduce（归约），是它们的主要思想，都是从函数式编程语言里借来的，还有从矢量编程语言里借来的特性。它极大地方便了编程人员在不会分布式并行编程的情况下，将自己的程序运行在分布式系统上。 当前的软件实现是指定一个Map（映射）函数，用来把一组键值对映射成一组新的键值对，指定并发的Reduce（归约）函数，用来保证所有映射的键值对中的每一个共享相同的键组。MapReduce是面向大数据并行处理的计算模型、框架和平台，它隐含了以下三层含义：1）MapReduce是一个基于集群的高性能并行计算平台（Cluster Infrastructure）。它允许用市场上普通的商用服务器构成一个包含数十、数百至数千个节点的分布和并行计算集群。2）MapReduce是一个并行计算与运行软件框架（Software Framework）。它提供了一个庞大但设计精良的并行计算软件框架，能自动完成计算任务的并行化处理，自动划分计算数据和计算任务，在集群节点上自动分配和执行任务以及收集计算结果，将数据分布存储、数据通信、容错处理等并行计算涉及到的很多系统底层的复杂细节交由系统负责处理，大大减少了软件开发人员的负担。3）MapReduce是一个并行程序设计模型与方法（Programming Model & Methodology）。它借助于函数式程序设计语言Lisp的设计思想，提供了一种简便的并行程序设计方法，用Map和Reduce两个函数编程实现基本的并行计算任务，提供了抽象的操作和并行编程接口，以简单方便地完成大规模数据的编程和计算处理Raft 是一种为了管理复制日志的一致性算法。它提供了和 Paxos 算法相同的功能和性能，但是它的算法结构和 Paxos 不同，使得 Raft 算法更加容易理解并且更容易构建实际的系统。为了提升可理解性，Raft 将一致性算法分解成了几个关键模块，例如leader人选举、日志复制和安全性。同时它通过实施一个更强的一致性来减少需要考虑的状态的数量。从一个用户研究的结果可以证明，对于学生而言，Raft 算法比 Paxos 算法更加容易学习。Raft 算法还包括一个新的机制来允许集群成员的动态改变，它利用重叠的大多数来保证安全性。 一致性算法允许一组机器像一个整体一样工作，即使其中一些机器出现故障也能够继续工作下去。正因为如此，一致性算法在构建可信赖的大规模软件系统中扮演着重要的角色。在过去的 10 年里，Paxos 算法统治着一致性算法这一领域：绝大多数的实现都是基于 Paxos 或者受其影响。同时 Paxos 也成为了教学领域里讲解一致性问题时的示例。 但是不幸的是，尽管有很多工作都在尝试降低它的复杂性，但是 Paxos 算法依然十分难以理解。并且，Paxos 自身的算法结构需要进行大幅的修改才能够应用到实际的系统中。这些都导致了工业界和学术界都对 Paxos 算法感到十分头疼。 和 Paxos 算法进行过努力之后，我们开始寻找一种新的一致性算法，可以为构建实际的系统和教学提供更好的基础。我们的做法是不寻常的，我们的首要目标是可理解性：我们是否可以在实际系统中定义一个一致性算法，并且能够比 Paxos 算法以一种更加容易的方式来学习。此外，我们希望该算法方便系统构建者的直觉的发展。不仅一个算法能够工作很重要，而且能够显而易见的知道为什么能工作也很重要。 Raft 一致性算法就是这些工作的结果。在设计 Raft 算法的时候，我们使用一些特别的技巧来提升它的可理解性，包括算法分解（Raft 主要被分成了leader人选举，日志复制和安全三个模块）和减少状态机的状态（相对于 Paxos，Raft 减少了非确定性和服务器互相处于非一致性的方式）。一份针对两所大学 43 个学生的研究表明 Raft 明显比 Paxos 算法更加容易理解。在这些学生同时学习了这两种算法之后，和 Paxos 比起来，其中 33 个学生能够回答有关于 Raft 的问题。 Raft 算法在许多方面和现有的一致性算法都很相似（主要是 Oki 和 Liskov 的 Viewstamped Replication），但是它也有一些独特的特性： 强leader：和其他一致性算法相比，Raft 使用一种更强的leader能力形式。比如，日志条目只从leader发送给其他的服务器。这种方式简化了对复制日志的管理并且使得 Raft 算法更加易于理解。leader选举：Raft 算法使用一个随机计时器来选举leader。这种方式只是在任何一致性算法都必须实现的心跳机制上增加了一点机制。在解决冲突的时候会更加简单快捷。成员关系调整：Raft 使用一种共同一致的方法来处理集群成员变换的问
NVIDIA CUDA 高度并行处理器编程（四）：CUDA性能优化
NVIDIA CUDA 高度并行处理器编程（四）：CUDA性能优化1. WARP 和线程执行 kernel 函数的执行速度在不同设备上差异很大，这是因为不同设备的资源约束不同。下面将会讨论资源约束的主要类型，以及它们如何限制 kernel 函数的执行性能。在特定设备上，可通过用一种资源代替另一种资源的方式来提高程序的性能。 1. WARP 和线程执行 （二）中写到，CUDA 设备将若干线程捆绑执行。把每个线程块划分成 warp。warp 的执行由 SIMD 硬件完成。线程块根据线程索引划分为 warp。当前
Spark中的键值对操作
1.PairRDD介绍 Spark为包含键值对类型的RDD提供了一些专有的操作。这些RDD被称为PairRDD。PairRDD提供了并行操作各个键或跨节点重新进行数据分组的操作接口。例如，PairRDD提供了reduceByKey()方法，可以分别规约每个键对应的数据，还有join()方法，可以把两个RDD中键相同的元素组合在一起，合并为一个RDD。2.创建Pair RDD 程序示例：...
Spark之RDD
目录 概述 一、RDD创建 Parallelized Collections External Datasets 二、RDD操作 Transformations Actions 三、RDD持久化 概述 在较高级别上，每个Spark应用程序都包含一个驱动程序，该程序运行在用户main函数中并且在集群上执行各种并行操作。Spark提供最主要的抽象是弹性分布式数据集(RDD)，它是跨...
《Python零基础到精通》学习笔记：3.6 经典函数实例
目录 3.6 经典函数实例 3.6.1 递归函数-汉诺塔的魅力 3.6.2 lambda函数-匿名函数的使用 3.6.3 Map-Reduce-映射与归约的思想 3.6 经典函数实例 3.6.1 递归函数-汉诺塔的魅力 在 Python 函数内部，我们可以去调用其他函数。所以如果一个函数在内部调用自身，这个函数我们就称为递归函数。 汉诺塔问题源于印度一个古老传说。相传大梵天创造世界的时候做了三根金刚石柱子，在一根柱子上从下往上按照大小顺序摞着64片黄金圆盘。大梵天命令婆罗门把圆盘从下面开
CUDA

580

社区成员

2,919

社区内容

发帖
与我相关
我的任务
社区描述
CUDA™是一种由NVIDIA推出的通用并行计算架构，该架构使GPU能够解决复杂的计算问题。 它包含了CUDA指令集架构（ISA）以及GPU内部的并行计算引擎。
社区管理员
  • CUDA编程社区
加入社区
  • 近7日
  • 近30日
  • 至今

加载中

查看更多榜单
社区公告
暂无公告

试试用AI创作助手写篇文章吧

+ 用AI写文章