0

我还是 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 没有改变输入数据），数据保持不变（好像我没有执行任何东西......位置没有调整）。

4

1 回答 1

0

好吧，我已经解决了问题..我在第三个内核本地组大小上做错了，在修复了所有事情都正确之后，对此感到抱歉

于 2012-12-30T19:26:41.893 回答