我有一个关于CUDA同步的问题.特别是,我需要一些澄清if语句的同步.我的意思是,如果我把__syncthreads()放在一个if语句的范围之内,这个if语句是由块内的一小部分线程命中的,那会怎么样?我以为一些线程将保持“永远”等待不会击中同步点的其他线程.所以,我写了并执行了一些示例代码来检查:
__global__ void kernel(float* vett,int n)
{
    int index = blockIdx.x*blockDim.x + threadIdx.x;
    int gridSize = blockDim.x*gridDim.x;

    while( index < n )
    {   
        vett[index] = 2;
        if(threadIdx.x < 10)
        {
            vett[index] = 100;
            __syncthreads();
        }
        __syncthreads();

        index += gridSize;
    }
}

令人惊讶的是,我观察到输出是一个非常“正常”(64个元素,块大小32):

100 100 100 100 100 100 100 100 100 100 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2
100 100 100 100 100 100 100 100 100 100 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2 2

所以我稍微修改了我的代码,方法如下:

__global__ void kernel(float* vett,int n)
{
    int index = blockIdx.x*blockDim.x + threadIdx.x;
    int gridSize = blockDim.x*gridDim.x;

    while( index < n )
    {   
        vett[index] = 2;
        if(threadIdx.x < 10)
        {
            vett[index] = 100;
            __syncthreads();
        }
        __syncthreads();
            vett[index] = 3;
        __syncthreads();

        index += gridSize;
    }
}

其输出为:

3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 
3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3

再次,我错了:我认为if语句中的线程在修改向量的元素后将保持在等待状态,永远不会离开if范围.
那么你可以澄清发生了什么吗?在同步点之后获得的线程是否解锁阻塞等待线程?
如果需要重现我的情况,我用SDK 4.2使用了CUDA Toolkit 5.0 RC.非常感谢.

解决方法

总之,行为是不明确的.所以它有时可以做你想要的,或者它可能不会,或者(很可能)会挂起或崩溃你的内核.

如果你真的很好奇内部的内容,你需要记住,线程不能独立执行,而是一个warp(一组32个线程).

这当然会导致条件分支的问题,其中条件不会在整个warp中统一评估.通过执行这两个路径,一个接一个地解决这个问题,每个路径被禁用,不应该执行该路径. IIRC在现有硬件上首先采用分支,然后执行不采用分支的路径,但是这种行为是未定义的,因此不能保证.

路径的这种单独执行持续到某种程度,编译器可以确定其保证由两个单独的执行路径(“再融合点”或“同步点”)的所有线程达到.当第一代码路径的执行到达该点时,它被停止并且代替执行第二代码路径.当第二条路径到达同步点时,所有线程将再次启用,并且执行从那里均匀地继续.

如果在同步之前遇到另一个条件分支,情况会变得更加复杂.这个问题是通过一堆仍然需要执行的路径解决的(幸运的是,堆栈的增长是有限的,因为我们可以为一个warp提供最多32个不同的代码路径).

插入同步点的位置是不确定的,甚至在体系结构之间略有不同,所以再也没有保证.从Nvidia获得的唯一(非官方)评论是,编译器非常适合找到最佳的同步点.然而,常常存在微妙的问题,可能会使您的最优点进一步下降,尤其是如果线程提前退出.

现在要了解__syncthreads()指令的行为(它转换成PTX中的一个bar.sync指令),重要的是要意识到这个指令不是每个线程都执行的,而是一次完整的转换(不管是否有任何的)线程被禁用或不被禁用),因为只有块的经线需要同步. warp的线程已经同步执行,并且进一步的同步将无效果(如果所有线程都被使能),或者当尝试从不同的条件代码路径同步线程时,会导致死锁.

您可以从此描述中了解您的特定代码行为.但请记住,所有这些都是未定义的,没有保证,依赖于具体行为可能会随时破坏您的代码.

您可能需要查看PTX manual的更多细节,特别是对于__syncthreads()编译的bar.sync指令.黄熙来的“Demystifying GPU Microarchitecture through Microbenchmarking” paper,以下由艾哈迈德参考,也值得一读.即使现在过时的架构和CUDA版本,有关条件分支和__syncthreads()的部分似乎仍然普遍有效.

CUDA:__syncthreads()里面的if语句的更多相关文章

  1. ios – 何时使用Semaphore而不是Dispatch Group?

    我会假设我知道如何使用DispatchGroup,为了解问题,我尝试过:结果–预期–是:为了使用信号量,我实现了:并在viewDidLoad方法中调用它.结果是:从概念上讲,dispachGroup和Semaphore都有同样的目的.老实说,我不熟悉:什么时候使用信号量,尤其是在与dispachGroup合作时–可能–处理问题.我错过了什么部分?

  2. ios – 当我的主线程阻塞时,如何获得断点/日志/增加的可见性?

    在对UI响应的永无止境的追求中,我想更多地了解主线程执行阻止操作的情况.我正在寻找某种“调试模式”或额外的代码,或钩子,或任何东西,从而我可以设置一个断点/日志/将被击中的东西,并允许我检查如果我的主要线程“自愿”用于I/O的块(或任何原因,真的),除了在循环结束时空闲.在过去,我已经使用循环观察器观察了跑步循环的时钟周期,这对于查看问题很有价值,但是在你可以检查的时候,为了做一个好主意,为时已晚

  3. ios – 如何在不阻塞主线程的情况下添加SCNNode?

    解决方法我不认为使用dispatchQueue可以解决这个问题.如果我替换其他任务而不是创建SCNNode它按预期工作,所以我认为问题与SceneKit有关.thisquestion的答案表明SceneKit有自己的私有后台线程,它将所有更改批量化.因此,无论我使用什么线程来创建我的SCNNode,它们都会在与渲染循环相同的线程中的同一队列中结束.我正在使用的丑陋的解决方法是在SceneKit的委托渲染器方法中一次添加一些节点,直到它们全部完成.

  4. 在Swift中应用Grand Central Dispatch(上

    在这两篇教程中,你会学到GCD的来龙去脉。起步libdispatch是Apple所提供的在IOS和OSX上进行并发编程的库,而GCD正是它市场化的名字。Swift中的闭包和OC中的块类似甚至于他们几乎就是可交换使用的。但OC中的块可以安全的替换成Swift中的闭包。再一次,这完全取决于GCD。QoS等级表示了提交任务的意图,使得GCD可以决定如何制定优先级。QOS_CLASS_USER_INteraCTIVE:userinteractive等级表示任务需要被立即执行以提供好的用户体验。

  5. 在Swift中应用Grand Central Dispatch 下

    通过使用dispatch_barrrier和dispatch_sync,你做到了让PhotoManager单例在读写照片时是线程安全的。还有,使用dispatch_async异步执行cpu密集型任务,从而为视图控制器初始化过程减负。幸运的是,dispatchgroups就是专为监视多个异步任务的完成情况而设计的。调度组调度组在一组任务都完成后会发出通知。在组内所有事件都完成时,GCDAPI提供了两种方式发送通知。打开PhotoManager.swift,替换downloadPhotosWithComple

  6. swift详解之十六-----------GCD基础部分

    当你了解了调度队列如何为你自己代码的不同部分提供线程安全后,GCD的优点就是显而易见的。这完全取决于GCD。这个队列就是用于发生消息给UIView或发送通知的。GCD的“艺术”归结为选择合适的队列来调度函数以提交你的工作。

  7. Realm Swift

    一旦带有主键的对象被添加到Realm之后,该对象的主键将不可修改。IgnoredProperties重写Object.ignoredProperties()可以防止Realm存储数据模型的某个属性。Realm将不会干涉这些属性的常规操作,它们将由成员变量提供支持,并且您能够轻易重写它们的setter和getter。所有的查询在Realm中都是延迟加载的,只有当属性被访问时,才能够读取相应的数据。

  8. 同步和异步

    如果是同步操作,它会阻塞当前线程并等待Block中的任务执行完毕,然后当前线程才会继续往下运行。并行队列中的任务根据同步或异步有不同的执行方式。同步执行异步执行串行队列当前线程,一个一个执行其他线程,一个一个执行并行队列当前线程,一个一个执行开很多线程,一起执行创建队列:主队列:这是一个特殊的串行队列。传入disPATCH_QUEUE_CONCURRENT表示创建并行队列。

  9. Swift--&gt;GCD,NSThread,NSBlockOperation多线程使用(主线程回调)

    应用程序开发,少不了的多线程,与多线程相关的就是线程同步.本文介绍Swift最简单的多线程使用.推荐阅读:http://www.jianshu.com/p/0b0d9b1f1f19看例子:1:获取线程基本的信息2:子线程的创建方法3:GCD(GrandCentraldispatch)队列的使用4:自定义queue5:NSBlockOperation和NSOperationQueue的使用6:子线程

  10. 完整详解swift GCD系列四dispatch_semaphore信号量

    viewmode=contents一何为信号量?简单来说就是控制访问资源的数量,比如系统有两个资源可以被利用,同时有三个线程要访问,只能允许两个线程访问,第三个应当等待资源被释放后再访问。其中value为信号量的初值,如果小于0则会返回NULL提高信号量copy

随机推荐

  1. 从C到C#的zlib(如何将byte []转换为流并将流转换为byte [])

    我的任务是使用zlib解压缩数据包(已接收),然后使用算法从数据中生成图片好消息是我在C中有代码,但任务是在C#中完成C我正在尝试使用zlib.NET,但所有演示都有该代码进行解压缩(C#)我的问题:我不想在解压缩后保存文件,因为我必须使用C代码中显示的算法.如何将byte[]数组转换为类似于C#zlib代码中的流来解压缩数据然后如何将流转换回字节数组?

  2. 为什么C标准使用不确定的变量未定义?

    垃圾价值存储在哪里,为什么目的?解决方法由于效率原因,C选择不将变量初始化为某些自动值.为了初始化这些数据,必须添加指令.以下是一个例子:产生:虽然这段代码:产生:你可以看到,一个完整的额外的指令用来移动1到x.这对于嵌入式系统来说至关重要.

  3. 如何使用命名管道从c调用WCF方法?

    更新:通过协议here,我无法弄清楚未知的信封记录.我在网上找不到任何例子.原版的:我有以下WCF服务我输出添加5行,所以我知道服务器是否处理了请求与否.我有一个.NET客户端,我曾经测试这一切,一切正常工作预期.现在我想为这个做一个非托管的C客户端.我想出了如何得到管道的名称,并写信给它.我从here下载了协议我可以写信给管道,但我看不懂.每当我尝试读取它,我得到一个ERROR_broKEN_P

  4. “这”是否保证指向C中的对象的开始?

    我想使用fwrite将一个对象写入顺序文件.班级就像当我将一个对象写入文件时.我正在游荡,我可以使用fwrite(this,sizeof(int),2,fo)写入前两个整数.问题是:这是否保证指向对象数据的开始,即使对象的最开始可能存在虚拟表.所以上面的操作是安全的.解决方法这提供了对象的地址,这不一定是第一个成员的地址.唯一的例外是所谓的标准布局类型.从C11标准:(9.2/20)Apointe

  5. c – 编译单元之间共享的全局const对象

    当我声明并初始化一个const对象时.两个cpp文件包含此标头.和当我构建解决方案时,没有链接错误,你会得到什么如果g_Const是一个非const基本类型!PrintInUnit1()和PrintInUnit2()表明在两个编译单元中有两个独立的“g_Const”具有不同的地址,为什么?

  6. 什么是C名称查找在这里? (&amp;GCC对吗?)

    为什么在第三个变体找到func,但是在实例化的时候,原始变体中不合格查找找不到func?解决方法一般规则是,任何不在模板定义上下文中的内容只能通过ADL来获取.换句话说,正常的不合格查找仅在模板定义上下文中执行.因为在定义中间语句时没有声明func,并且func不在与ns::type相关联的命名空间中,所以代码形式不正确.

  7. c – 在输出参数中使用auto

    有没有办法在这种情况下使用auto关键字:当然,不可能知道什么类型的.因此,解决方案应该是以某种方式将它们合并为一个句子.这可用吗?解决方法看起来您希望默认初始化给定函数期望作为参数的类型的对象.您无法使用auto执行此操作,但您可以编写一个特征来提取函数所需的类型,然后使用它来声明您的变量:然后你就像这样使用它:当然,只要你重载函数,这一切都会失败.

  8. 在C中说“推动一切浮动”的确定性方式

    鉴于我更喜欢将程序中的数字保留为int或任何内容,那么使用这些数字的浮点数等效的任意算术最方便的方法是什么?说,我有我想写通过将转换放在解析的运算符树叶中,无需将表达式转化为混乱是否可以使用C风格的宏?应该用新的类和重载操作符完成吗?解决方法这是一个非常复杂的表达.更好地给它一个名字:现在当您使用整数参数调用它时,由于参数的类型为double,因此使用常规的算术转换将参数转换为double用C11lambda……

  9. objective-c – 如何获取未知大小的NSArray的第一个X元素?

    在objectiveC中,我有一个NSArray,我们称之为NSArray*largeArray,我想要获得一个新的NSArray*smallArray,只有第一个x对象…

  10. c – Setprecision是混乱

    我只是想问一下setprecision,因为我有点困惑.这里是代码:其中x=以下:方程的左边是x的值.1.105=1.10应为1.111.115=1.11应为1.121.125=1.12应为1.131.135=1.14是正确的1.145=1.15也正确但如果x是:2.115=2.12是正确的2.125=2.12应为2.13所以为什么在一定的价值是正确的,但有时是错误的?请启发我谢谢解决方法没有理由期望使用浮点系统可以正确地表示您的帖子中的任何常量.因此,一旦将它们存储在一个双变量中,那么你所拥有的确切的一

返回
顶部