问题描述
|
考虑以下代码:
__global__ void kernel(int *something) {
extern __shared__ int shared_array[];
// Some operations on shared_array here.
}
是否可以将整个shared_array设置为某个值-例如0-是否未明确寻址某个线程中的每个单元?
解决方法
否。共享内存未初始化。您必须以某种方式自行初始化它...
摘自《 CUDA C编程指南3.2》第B.2.4.2节第2段:
__shared__
变量不能在其声明中包含初始化。
这也会丢弃共享变量的非平凡的默认构造函数。
, 您可以像这样高效地并行初始化共享数组
// if SHARED_SIZE == blockDim.x,eliminate this loop
for (int i = threadIdx.x; i < SHARED_SIZE; i += blockDim.x)
shared_array[i] = INITIAL_VALUE;
__syncthreads();
, 是的你可以。您可以指定块中的第一个线程对其进行设置,而另一个则不这样:
extern __shared__ unsigned int local_bin[]; // Size specified in kernel call
if (threadIdx.x == 0) // Wipe on first thread - include \" && threadIdx.y == 0\" and \" && threadIdx.z == 0\" if threadblock has 2 or 3 dimensions instead of 1.
{
// For-loop to set all local_bin array indexes to specified value here - note you cannot use cudaMemset as it translates to a kernel call itself
}
// Do stuff unrelated to local_bin here
__syncthreads(); // To make sure the memset above has completed before other threads start writing values to local_bin.
// Do stuff to local_bin here
理想情况下,您应该在syncthreads调用之前做尽可能多的工作,因为这允许所有其他线程在memset完成之前完成其工作-显然,这仅在工作可能具有完全不同的线程完成时间的情况下才重要,例如,如果存在条件分支。
请注意,对于线程0 \“ setting \” for循环,您需要将local_bin数组的大小作为参数传递给内核,以便您知道要迭代的数组的大小。
原始概念源