我还是 OpenCL 的新手，我正在用 Nvidia 示例进行一些测试，整个程序由 5 个内核组成，这些内核按顺序执行（1、2、3、4、5）。
第一个内核简单地获取位置数据、速度数据，并应用重力和基本碰撞检测，然后调整该位置和速度……这个内核完美运行，没有问题。
这是第一个内核：
__kernel void integrate(
__global float4 *d_Pos, //input/output
__global float4 *d_Vel, //input/output
__constant simParams_t *params,
float deltaTime,
uint numParticles
){
const uint index = get_global_id(0);
if(index >= numParticles)
return;
float4 pos = d_Pos[index];
float4 vel = d_Vel[index];
pos.w = 1.0f;
vel.w = 0.0f;
//Gravity
vel += (float4)(params->gravity.x, params->gravity.y, params->gravity.z, 0) * deltaTime;
vel *= params->globalDamping;
//Advance pos
pos += vel * deltaTime;
//Collide with cube
if(pos.x < -1.0f + params->particleRadius){
pos.x = -1.0f + params->particleRadius;
vel.x *= params->boundaryDamping;
}
if(pos.x > 1.0f - params->particleRadius){
pos.x = 1.0f - params->particleRadius;
vel.x *= params->boundaryDamping;
}
if(pos.y < -1.0f + params->particleRadius){
pos.y = -1.0f + params->particleRadius;
vel.y *= params->boundaryDamping;
}
if(pos.y > 1.0f - params->particleRadius){
pos.y = 1.0f - params->particleRadius;
vel.y *= params->boundaryDamping;
}
if(pos.z < -1.0f + params->particleRadius){
pos.z = -1.0f + params->particleRadius;
vel.z *= params->boundaryDamping;
}
if(pos.z > 1.0f - params->particleRadius){
pos.z = 1.0f - params->particleRadius;
vel.z *= params->boundaryDamping;
}
//Store new position and velocity
d_Pos[index] = pos;
d_Vel[index] = vel;
}
第二个内核将这些位置作为输入并输出另一种数据（一些索引），但它不会更改位置数据。
第三个内核正在对第二个内核输出进行调整（从不涉及位置数据的第二个内核获取数据）。
现在解决问题...第四个内核；这需要位置数据和速度数据（来自第一个内核），从第三个内核获取调整后的数据，输出另一个位置和速度数据（这些位置和速度的指针完全不同）
这是第四个内核：
__kernel void findCellBoundsAndReorder(
__global uint *d_CellStart, //output: cell start index
__global uint *d_CellEnd, //output: cell end index
__global float4 *d_ReorderedPos, //output: reordered by cell hash positions
__global float4 *d_ReorderedVel, //output: reordered by cell hash velocities
__global const uint *d_Hash, //input: sorted grid hashes
__global const uint *d_Index, //input: particle indices sorted by hash
__global const float4 *d_Pos, //input: positions array sorted by hash
__global const float4 *d_Vel, //input: velocity array sorted by hash
__local uint *localHash, //get_group_size(0) + 1 elements
uint numParticles
){
uint hash;
const uint index = get_global_id(0);
//Handle case when no. of particles not multiple of block size
if(index < numParticles){
hash = d_Hash[index];
//Load hash data into local memory so that we can look
//at neighboring particle's hash value without loading
//two hash values per thread
localHash[get_local_id(0) + 1] = hash;
//First thread in block must load neighbor particle hash
if(index > 0 && get_local_id(0) == 0)
localHash[0] = d_Hash[index - 1];
}
barrier(CLK_LOCAL_MEM_FENCE);
if(index < numParticles){
//Border case
if(index == 0)
d_CellStart[hash] = 0;
//Main case
else{
if(hash != localHash[get_local_id(0)])
d_CellEnd[localHash[get_local_id(0)]] = d_CellStart[hash] = index;
};
//Another border case
if(index == numParticles - 1)
d_CellEnd[hash] = numParticles;
//Now use the sorted index to reorder the pos and vel arrays
uint sortedIndex = d_Index[index];
float4 pos = d_Pos[sortedIndex];
float4 vel = d_Vel[sortedIndex];
d_ReorderedPos[index] = pos;
d_ReorderedVel[index] = vel;
}
}
问题是如果我单独执行内核 1（或 1+2，或 1+2+3）位置和速度从第一个内核正确调整。
但是如果我执行内核 1+2+3+4（尽管内核 4 没有改变输入数据），数据保持不变（好像我没有执行任何东西......位置没有调整）。