我现在用100个threads读取一个10*10的数组. 如果数组中所有的元素都为0, 那么不执行任何操作; 若有至少一个非0, 那执行一些既定的操作. 即: 我需要用一个标志Flag来指示该数组中是否有非0元素.
shared int Flag=0;
若某一thread读到的数据非0, 则置Flag=1;
__syncthreads();
但这样一来,就相当于同一个block 中的可能所有threads几乎会在同一时间对Flag进行写操作 (假设数组中所有的元素都不为0),在程序中可以这样用吗? 那这样会不会出现竞争冒险呢?
感谢!
LZ您好：
一个warp的线程对shared memory中同一位置写入的时候，仅有一个线程能写入成功，且不保证是哪个线程成功。所以您这样做是有风险的，请勿这样。
祝您编码顺利~
经玫瑰斑竹提示，上述的说法有考虑不周之处，请以下面修改的为准。
1：手册的确表明，“一个warp的线程对shared memory中同一位置写入的时候，仅有一个线程能写入成功，且不保证是哪个线程成功。”但是您可以安排读到0的线程什么也不干，读到非0的线程去写flag，此时如果warp内有一个线程读到非0的值，那么直接是可以写成功的；如果有多于一个线程读到非0的值，那么这些线程中会有一个写成功，依然能完成置flag。
（我之前习惯性地以为读到0的线程也要去置flag，实际上不需要。）
2：以及，您需要在开始读取数据并置flag之前，先选定一个线程对flag初始化为0，并__syncthreads().
大致如此，之前考虑不周造成误导，深表歉意。
以及，您可以换用其他的实现方式，比如采用原子操作对shared memory中的flag进行标记或累加；或者您可以对所有数据进行规约操作，统计出有没有非零数据，最后再置标记位，等等。
非常感谢ICE版主!
您讲得很清楚,受教了:)
您过奖了，应该感谢玫瑰斑竹提醒，否则就误导你了。
欢迎您常来论坛~
版主您好, 感谢您的回复.我刚才又想到一个问题: 如果一个WARP 里的thread都去读同一个地址呢? 这样每个thread都能得到正确的结果吗?
另外, 用__global__声明的函数中, 可以去调用别的函数吗? 如果可以, 参数传递是否与传统C语言一样呢?
谢谢啊!
LZ您好：
1：读取无妨，可以通过广播机制使得所有的读取线程获得正确的结果。
以及，需要注意的是，请不要混合普通读取和原子写入。
2：__global__函数可以调用__device__函数，参数传递和C中一样。
此外在SM3.5的硬件上可以调用__global__函数，这是其特有的动态并行功能。
大致如此，祝您好运~
:handshake 
非常感谢!
不客气，欢迎您常来论坛~
版主:不好意思我还有一个问题…
用__global__定义的内核函数里,如果用通常的方法来定义变量(如float A, int B), 程序在执行时是否每个 thread 都会在物理空间里为自己分配对应的变量呢?
以上为例,如果一个block 里有 1024 threads, 程序运行时将会在L1里分配 (sizeof(A) + sizeof(B))*1024 bytes 的空间呢? 而每个block最大的L1是48KB, 也就是说在内核函数里定义的变量(不包括shared)最多不能超过 48 bytes, 对吗?
谢谢啦…
LZ您好：
1：kernel里面定义的变量是每个线程私有的，也就是每个线程有自己的副本，不能跨线程访问。
2：这些变量一般会使用寄存器资源，当寄存器资源不够用的时候，还可以使用local memory。这个和L1 cache无直接关系。所以您的说法是不正确的。
祝您编码顺利~
:handshake
多谢版主的解答:)
不客气的，欢迎您常来论坛~