-1

3D ポイント クラウドがあり、ピクセルをイメージ プレーンに投影します。一部の 3D ポイントは同じピクセルにマッピングされるため、Z 値が最も低いピクセルのみをカメラに表示したいと考えています。深度値を追跡するために、Z バッファー (float 配列) を使用します。ここにいくつかの擬似コードがあります:

// Initialize z-Buffer with max depth (99999.9f)
// Go through every point of point cloud...
// Project Point (x,y,z) to image plane (u,v) 
int newIndex = v*imgWidth+u;
float oldDepth = zbuffer[newIndex];

if (z < oldDepth){
  zbuffer[newIndex] = z; // put z value in buffer
  outputImg[newIndex] = pointColor[i]; // put pixel in resulting image
}

これの完全に機能するシングルコア CPU バージョンがあります。

cuda バージョンは問題なく見え、非常に高速ですが、z-test が役割を果たす領域だけが非常に「縞模様」になっています。これは、一部の背景ポイントが前景ピクセルを上書きしていることを意味すると思います。また、カラー画像を見ると、画像には存在しない色のランダムな色の筋が見られます。

CUDA バージョンは次のようになります。

//Initialize, kernel, project, new coordinates...

const float oldDepth = outputDepth[v * outputMaxWidth + u];

if (z < oldDepth){
  outputDepth[v * outputMaxWidth + u] = z;
  const int inputColorIndex = yIndex * inputImageStep + 3*xIndex;
  const int outputColorIndex = yIndex * outputImageStep + 3*xIndex;
  outputImage[outputColorIndex] = inputImage[inputColorIndex]; //B
  outputImage[outputColorIndex + 1] = inputImage[inputColorIndex + 1]; //G
  outputImage[outputColorIndex + 2] = inputImage[inputColorIndex + 2]; //R
}

ここで並行性が問題になると思います。1 つのスレッドがこのピクセルの最も近い Z 値を Z バッファーに書き込むことがありますが、同時に別のスレッドが古い値を読み取り、正しい値を上書きします。

CUDAでこれを防ぐにはどうすればよいですか?

Edit1: ブロック サイズを (16,16) から (1,1) に縮小すると、縞模様は少なくなりますが、1 ピクセルの穴のように見えます。

Edit2:これは最小限の例です:

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

cudaError_t insertToZBuffer(int *z, const int *a, unsigned int size);

__global__ void zbufferKernel(int *z, const int *a)
{
    int i = threadIdx.x;
    if (a[i] < z[0]){
        z[0] = a[i]; //  all mapped to pixel index 0        
    }    
}

int main(){
    for (int i = 0; i < 20; ++i){
        const int arraySize = 5;
        const int a[arraySize] = { 1, 7, 3, 40, 5 }; // some depth values which get mapped all to index 0
        int z[arraySize] = { 999 }; // large depth value

        insertToZBuffer(z, a, arraySize);

        printf("{%d,%d,%d,%d,%d}\n", z[0], z[1], z[2], z[3], z[4]);
        cudaDeviceReset();

    }   
    return 0;
}

cudaError_t insertToZBuffer(int *z, const int *a, unsigned int size){
    int *dev_a = 0;
    int *dev_z = 0;
    cudaError_t cudaStatus;
    cudaStatus = cudaSetDevice(0);
    cudaStatus = cudaMalloc((void**)&dev_z, size * sizeof(int));
    cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
    cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
    cudaStatus = cudaMemcpy(dev_z, z, size * sizeof(int), cudaMemcpyHostToDevice);
    zbufferKernel<<<1, size >>>(dev_z, dev_a);
    cudaStatus = cudaGetLastError();
    cudaStatus = cudaDeviceSynchronize();   
    cudaStatus = cudaMemcpy(z, dev_z, size * sizeof(int), cudaMemcpyDeviceToHost);

    cudaFree(dev_z);
    cudaFree(dev_a);

    return cudaStatus;
}

インデックス 0 の z の値は、それが最小値であるため 1 である必要がありますが、a の最後の値である 5 です。

4

1 に答える 1

-1

コメントのおかげで解決した方法は次のとおりです。

z 値が小さい場合は、atomicCAS (float を int にキャスト) を使用して z バッファーに書き込みました。現在のスレッドの z 値が大きい場合は、単純に戻ります。最後に、バッファに書き込みを行った残りのすべてのスレッド (__syncthreads()) を同期し、それらの z 値が最後の値であるかどうかを確認します。それが本当なら、この位置のピクセル値にポイント カラーを書き込みます。

編集：atomicMinだけを使用する必要がありました...

于 2017-01-04T11:36:14.623 に答える