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;
}

2 番目のカーネルは、これらの位置を入力として受け取り、別の種類のデータ (いくつかのインデックス) を出力しますが、位置データは変更しません。

3 番目のカーネルは、2 番目のカーネル出力の調整を行っています (位置データに触れない 2 番目のカーネルからデータを取得します)。

さて、問題は...4番目のカーネルです。これは位置データと速度データを (最初のカーネルから) 取得し、調整されたデータを 3 番目のカーネルから取得し、別の位置と速度データを出力します (これらの位置と速度のまったく異なるポインター)。

4 番目のカーネルは次のとおりです。

__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

さて、私は問題を理解しました..すべてが正しくなったことを修正した後、3番目のカーネルローカルグループサイズで間違いを犯していました、申し訳ありません

于 2012-12-30T19:26:41.893 に答える