我有一个关于CUDA同步的问题.特别是,我需要对if语句中的同步进行一些澄清.我的意思是,如果我将__syncthreads()放在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 …我正在开发一些具有可变数量线程的C++多核程序,我想知道如何设置一个合适的(实际上是"最好的")亲和力.我使用Boost-threads,所以我可以调用get_hardware_concurrency()来了解有多少逻辑内核.到目前为止,我写了一个映射"第n个线程到第n个逻辑核心",但由于多插槽处理器和超线程,它并不是最聪明的事情.我的程序总是像SIMD一样,所以线程之间没有任何共享,如果是HT计算机,我想以我能想象的最聪明的方式将线程绑定到逻辑核心:第一个物理上的第一个逻辑核心,第2个物理上的第1个逻辑,...,第1个物理上的第1个逻辑,第1个物理上的第2个逻辑,依此类推.
我发现了很多内容,讨论了如何发现HT是否启用(CPUID)以及如何确定逻辑和物理内核PER包.我知道我必须处理一些汇编代码,它并没有吓到我,但我真的找不到如何知道有关逻辑内核,物理内核和软件包的完整信息以及操作系统如何处理所有这些信息.
作为最简洁的我可以:我怎么知道OS(Windows和Linux)引用的线程的确切位置(物理核心和包)为N-th？
几天前,我正在比较我的一些代码的性能,我执行一个非常简单的替换和Thrust实现相同的算法.我发现一个数量级(!)的不匹配有利于Thrust,所以我开始让我的调试器"冲浪"到他们的代码中以发现魔法发生的地方.
令人惊讶的是,我发现我的非常直接的实现实际上非常类似于他们的,一旦我摆脱了所有的functor东西并得到了细节.我看到Thrust有一个聪明的方法来决定块_size和grid_size(顺便说一句:确切地说,它是如何工作的？!),所以我只是采取了他们的设置并再次执行我的代码,因为它们非常相似.我增加了几微秒,但情况几乎相同.然后,最后,我不知道为什么,只是"尝试"我删除了我的内核和BINGO之后的cudaThreadSynchronize()!我将差距归零(并且更好),并获得了整个执行时间的数量级.访问我的数组值我发现它们完全符合我的预期,所以正确执行.
现在的问题是:我什么时候可以摆脱cudaThreadSynchronize(et similia)？为什么会造成如此巨大的开销？我看到Thrust本身在最后没有同步(synchronize_if_enabled(const char*message)是一个NOP,如果没有定义宏__THRUST_SYNCHRONOUS而且它不是).细节和代码如下.
// my replace code
template <typename T>
__global__ void replaceSimple(T* dev, const int n, const T oldval, const T newval)
{
const int gridSize = blockDim.x * gridDim.x;
int index = blockIdx.x * blockDim.x + threadIdx.x;
while(index < n)
{
if(dev[index] == oldval)
dev[index] = newval;
index += gridSize;
}
}
// replace invocation - not in main because of cpp - cu separation
template <typename T>
void callReplaceSimple(T* dev, const int n, const T …