__shfl_down和__shfl_down_sync给出不同的结果

作者：无痕 | 来源：互联网 | 2022-12-15 13:28

如何解决《__shfl_down和__shfl_down_sync给出不同的结果》经验，为你挑选了1个好方法。

我对__shfl_down和__shfl_down_sync感到困惑,他们给出了不同的结果.

__global__ void shufledown1(double* a, double *b,double *c, int N)
{
    double  temp = 2.0;
    __syncthreads();

   for (int offset = 32/2; offset > 0; offset /= 2){
         temp+=__shfl_down(temp, offset,32);
   }
    printf("%d %f %d \n",threadIdx.x ,temp,blockDim.x * gridDim.x);
}
__global__ void shufledown2(double* a, double *b,double *c, int N)
{
    double  temp = 2.0;
    __syncthreads();

   for (int offset = 32/2; offset > 0; offset /= 2){
       temp+=__shfl_down_sync(temp, offset,32)
   }
    printf("%d %f %d \n",threadIdx.x ,temp,blockDim.x * gridDim.x);
}

第一个给出:

0 64.000000 64 
'''''
''''
''''
63 64.000000 64

第二个给出:

0 33.000000 64
'''''
''''
''''
63 33.000000 64

内核使用1个块64个线程运行.问候

1> Robert Crove..：

除了一个不同的名称之外,_syncwarp shuffle函数的版本也有不同的原型,如文档中所示.第一个参数是掩码参数.

您似乎试图以相同的方式使用这两个函数:

     temp+=__shfl_down(temp, offset,32);

temp+=__shfl_down_sync(temp, offset,32);

但这是不正确的.要_sync以类似的方式使用该版本,您应该:

temp+=__shfl_down_sync(0xFFFFFFFF, temp, offset,32);

当我进行更改时,您的代码正确地为我运行:

#include 

__global__ void shufledown1(double* a, double *b,double *c, int N)
{
    double  temp = 2.0;
    __syncthreads();

   for (int offset = 32/2; offset > 0; offset /= 2){
         temp+=__shfl_down(temp, offset,32);
   }
    printf("%d %f %d \n",threadIdx.x ,temp,blockDim.x * gridDim.x);
}
__global__ void shufledown2(double* a, double *b,double *c, int N)
{
    double  temp = 2.0;
    __syncthreads();

   for (int offset = 32/2; offset > 0; offset /= 2){
       temp+=__shfl_down_sync(0xFFFFFFFF, temp, offset,32);
   }
    printf("%d %f %d \n",threadIdx.x ,temp,blockDim.x * gridDim.x);
}


int main(){
    double *a = NULL, *b = NULL, *c = NULL;
    shufledown1<<<1,64>>>(a, b, c, 0);
    cudaDeviceSynchronize();
    shufledown2<<<1,64>>>(a, b, c, 0);
    cudaDeviceSynchronize();
}
$ nvcc -arch=sm_60 -o t1358 t1358.cu
t1358.cu(9): warning: function "__shfl_down(double, unsigned int, int)"
/usr/local/cuda/bin/..//include/sm_30_intrinsics.hpp(453): here was declared deprecated ("__shfl_down() is deprecated in favor of __shfl_down_sync() and may be removed in a future release (Use -Wno-deprecated-declarations to suppress this warning).")

ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 49; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 52; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 63; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 66; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 77; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 80; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 91; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 94; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 105; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
ptxas /tmp/tmpxft_000045b6_00000000-5_t1358.ptx, line 108; warning : Instruction 'shfl' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version
$ ./t1358
0 64.000000 64
1 64.000000 64
2 64.000000 64
3 64.000000 64
4 64.000000 64
5 64.000000 64
6 64.000000 64
7 64.000000 64
8 64.000000 64
9 64.000000 64
10 64.000000 64
11 64.000000 64
12 64.000000 64
13 64.000000 64
14 64.000000 64
15 64.000000 64
16 64.000000 64
17 64.000000 64
18 64.000000 64
19 64.000000 64
20 64.000000 64
21 64.000000 64
22 64.000000 64
23 64.000000 64
24 64.000000 64
25 64.000000 64
26 64.000000 64
27 64.000000 64
28 64.000000 64
29 64.000000 64
30 64.000000 64
31 64.000000 64
32 64.000000 64
33 64.000000 64
34 64.000000 64
35 64.000000 64
36 64.000000 64
37 64.000000 64
38 64.000000 64
39 64.000000 64
40 64.000000 64
41 64.000000 64
42 64.000000 64
43 64.000000 64
44 64.000000 64
45 64.000000 64
46 64.000000 64
47 64.000000 64
48 64.000000 64
49 64.000000 64
50 64.000000 64
51 64.000000 64
52 64.000000 64
53 64.000000 64
54 64.000000 64
55 64.000000 64
56 64.000000 64
57 64.000000 64
58 64.000000 64
59 64.000000 64
60 64.000000 64
61 64.000000 64
62 64.000000 64
63 64.000000 64
0 64.000000 64
1 64.000000 64
2 64.000000 64
3 64.000000 64
4 64.000000 64
5 64.000000 64
6 64.000000 64
7 64.000000 64
8 64.000000 64
9 64.000000 64
10 64.000000 64
11 64.000000 64
12 64.000000 64
13 64.000000 64
14 64.000000 64
15 64.000000 64
16 64.000000 64
17 64.000000 64
18 64.000000 64
19 64.000000 64
20 64.000000 64
21 64.000000 64
22 64.000000 64
23 64.000000 64
24 64.000000 64
25 64.000000 64
26 64.000000 64
27 64.000000 64
28 64.000000 64
29 64.000000 64
30 64.000000 64
31 64.000000 64
32 64.000000 64
33 64.000000 64
34 64.000000 64
35 64.000000 64
36 64.000000 64
37 64.000000 64
38 64.000000 64
39 64.000000 64
40 64.000000 64
41 64.000000 64
42 64.000000 64
43 64.000000 64
44 64.000000 64
45 64.000000 64
46 64.000000 64
47 64.000000 64
48 64.000000 64
49 64.000000 64
50 64.000000 64
51 64.000000 64
52 64.000000 64
53 64.000000 64
54 64.000000 64
55 64.000000 64
56 64.000000 64
57 64.000000 64
58 64.000000 64
59 64.000000 64
60 64.000000 64
61 64.000000 64
62 64.000000 64
63 64.000000 64
$

对于新代码或新维护,您应该只使用这些_sync版本.

有关mask参数用法的更多示例,请参阅此博客

您是否阅读了我指出的文档？例如，本段：“新的* _sync shfl内部函数采用掩码表示参与调用的线程。必须为每个参与线程设置一个代表线程的通道ID的位，以确保它们在内部函数之前正确收敛由硬件执行。在mask中命名的所有未退出线程必须在相同的mask中执行相同的内在函数，否则结果是不确定的。”

推荐阅读

list
PhysioNet生理信号处理（三）WFDB Toolbox for Matlab的安装和使用方法

本文介绍了PhysioNet网站提供的生理信号处理工具箱WFDB Toolbox for Matlab的安装和使用方法。通过下载并添加到Matlab路径中或直接在Matlab中输入相关内容，即可完成安装。该工具箱提供了一系列函数，可以方便地处理生理信号数据。详细的安装和使用方法可以参考本文内容。 ... [详细]

蜡笔小新 2023-12-13 20:46:48
list
iOS Metal shader中的if条件判断限制及处理方式

本文讨论了在iOS平台中的Metal框架中，对于if语句中的判断条件的限制和处理方式。作者提到了在Metal shader中，判断条件不能写得太长太复杂，否则可能导致程序停留或没有响应。作者还分享了自己的经验，建议在CPU端进行处理，以避免出现问题。 ... [详细]

蜡笔小新 2023-12-12 19:01:06
java
JDK源码学习之HashTable(附带面试题)的学习笔记

本文介绍了JDK源码学习之HashTable(附带面试题)的学习笔记，包括HashTable的定义、数据类型、与HashMap的关系和区别。文章提供了干货，并附带了其他相关主题的学习笔记。 ... [详细]

蜡笔小新 2023-12-12 13:05:17
java
Learning to Paint with Model-based Deep Reinforcement Learning

本文介绍了一种基于模型的深度强化学习方法，通过结合神经渲染器，教机器像人类画家一样进行绘画。该方法能够生成笔画的坐标点、半径、透明度、颜色值等，以生成类似于给定目标图像的绘画。文章还讨论了该方法面临的挑战，包括绘制纹理丰富的图像等。通过对比实验的结果，作者证明了基于模型的深度强化学习方法相对于基于模型的DDPG和模型无关的DDPG方法的优势。该研究对于深度强化学习在绘画领域的应用具有重要意义。 ... [详细]

蜡笔小新 2023-12-11 10:27:44
java
go利用(*interface{})(nil)传递参数类型的原理及应用

本文介绍了在go语言中利用(*interface{})(nil)传递参数类型的原理及应用。通过分析Martini框架中的injector类型的声明，解释了values映射表的作用以及parent Injector的含义。同时，讨论了该技术在实际开发中的应用场景。 ... [详细]

蜡笔小新 2023-12-10 11:37:12
filter
如何在Jquery中通过其样式值获取元素

本文介绍了如何在Jquery中通过元素的样式值获取元素，并将其赋值给一个变量。提供了5种解决方案供参考。 ... [详细]

蜡笔小新 2023-12-10 04:03:11
filter
python 终止函数命令_如何使“停止”按钮终止已经在Tkinter（Python）中运行的“启动”函数...

我用Tkinter制作了一个图形用户界面，有两个主按钮：“开始”和“停止”。请您就如何使用“停止”按钮终止“开始”按钮为以下代码调用的已运行功能提供建议 ... [详细]

蜡笔小新 2023-10-17 20:02:38
input
在 IMDB 情感分类任务上训练双向 LSTM

代码如下在这里插入代码片from__future__importprint_functionimportnumpyasnpfromkeras.preprocessingimpo ... [详细]

蜡笔小新 2023-10-17 12:30:42
filter
Wordpress：仅显示未来的帖子减去一天 - Wordpress: Only show future posts minus one day

SoIhavealoopthatrunsperfectforeventsandonlyshowsfutureposts.TheissueisthatIwould ... [详细]

蜡笔小新 2023-10-17 09:11:29
java
Highcharts翻译系列之二十：曲线图例子（二）

Highcharts翻译系列之二十：曲线图例子（二）代码 ... [详细]

蜡笔小新 2023-10-16 17:30:36
java
delphiParallel之TTask初试

unitUnit1;interfaceusesWinapi.Windows,Winapi.Messages,System.SysUtils,System.Variants,Syst ... [详细]

蜡笔小新 2023-10-16 13:37:53
export
jquery中如何获取文件的大小

最近百度了好几种方式都没有效果针对这种a标签直接open的方式我也是醉了因为要对这个导出文件大小进行限制,当文件大于100mb的时候,就会有提示并且不让下载对于这种前端计算是很难 ... [详细]

蜡笔小新 2023-10-16 13:12:34
export
Python中的异常处理示例及其应用

Python异常处理python提供了两个非常重要的功能来处理python程序在运行中出现的异常和错误。你可以使用该功能来调试python程序。异常处理:本站Python教程会 ... [详细]

蜡笔小新 2023-10-14 20:22:40
export
./configure make error

error..libtool:line838:X--tagCC:commandnotfound..libtool:line871:libtool:ignoringunknown ... [详细]

蜡笔小新 2023-10-13 18:57:09
export
[Python] future线程处理并发的核心源码详解

ThreadPoolExecutor核心工作详解ThreadPoolExecutor内部是如何工作的？什么是future对象？线程什么时候被创建&#x ... [详细]

蜡笔小新 2023-10-13 13:35:18

无痕

这个家伙很懒，什么也没留下！

Tags | 热门标签

RankList | 热门文章