你能发布整个内核代码吗?我必须对参数和私有变量做出假设。
看起来组中有 nt 个工作项,ti 代表当前工作项。当循环执行时,组中的每个项目将只复制单个元素。通常,此副本来自全局数据源。第一个障碍强制工作项等待,直到其他项完成复制。这是必要的,因为组中的每个工作项都需要读取从其他每个工作项复制的数据。这些值不应该相同,因为每个工作项的 ti 应该不同。(但对于第一个循环,jb*nt 仍然为零)
这是整个内核代码:
__kernel
void
nbody_sim(
__global float4* pos ,
__global float4* vel,
int numBodies,
float deltaTime,
float epsSqr,
__local float4* localPos,
__global float4* newPosition,
__global float4* newVelocity)
{
unsigned int tid = get_local_id(0);
unsigned int gid = get_global_id(0);
unsigned int localSize = get_local_size(0);
// Number of tiles we need to iterate
unsigned int numTiles = numBodies / localSize;
// position of this work-item
float4 myPos = pos[gid];
float4 acc = (float4)(0.0f, 0.0f, 0.0f, 0.0f);
for(int i = 0; i < numTiles; ++i)
{
// load one tile into local memory
int idx = i * localSize + tid;
localPos[tid] = pos[idx];
// Synchronize to make sure data is available for processing
barrier(CLK_LOCAL_MEM_FENCE);
// calculate acceleration effect due to each body
// a[i->j] = m[j] * r[i->j] / (r^2 + epsSqr)^(3/2)
for(int j = 0; j < localSize; ++j)
{
// Calculate acceleartion caused by particle j on particle i
float4 r = localPos[j] - myPos;
float distSqr = r.x * r.x + r.y * r.y + r.z * r.z;
float invDist = 1.0f / sqrt(distSqr + epsSqr);
float invDistCube = invDist * invDist * invDist;
float s = localPos[j].w * invDistCube;
// accumulate effect of all particles
acc += s * r;
}
// Synchronize so that next tile can be loaded
barrier(CLK_LOCAL_MEM_FENCE);
}
float4 oldVel = vel[gid];
// updated position and velocity
float4 newPos = myPos + oldVel * deltaTime + acc * 0.5f * deltaTime * deltaTime;
newPos.w = myPos.w;
float4 newVel = oldVel + acc * deltaTime;
// write to global memory
newPosition[gid] = newPos;
newVelocity[gid] = newVel;
}
每个工作组都有“numTiles”工作组和“localSize”工作项。
“gid”是全局索引,“tid”是本地索引。
让我们从循环“for(int i = 0; i < numTiles; ++i)”的第一次迭代开始,其中“i=0”:
如果我举个例子:
numTiles = 4, localSize = 25 和 numBodies = 100 = 工作项数。
然后,在执行时,如果我有 gid = 80,则 tid = 5,idx = 5,第一个分配将是:localPos[5] = pos[5]
现在,我取 gid = 5,然后 tid = 5 和 idx = 5,我将有相同的分配:localPos[5] = pos[5]
所以,据我了解,在第一次迭代和第一个“障碍”之后,每个工作项都包含相同的本地数组“localPos”,即第一个全局块的子数组,即“pos[0: 24]”。
这是对发生的事情的一个很好的解释吗?