2

要素が多くの制約に準拠している場合、ベクトルから要素のインデックスを抽出するために、Thrust のストリーム圧縮機能 (copy_if) を使用したいと考えています。これらの制約の 1 つは、隣接する要素の値に依存します (2D では 8、3D では 26)。私の質問は、Thrust で要素の隣人を取得するにはどうすればよいですか?

「copy_if」のファンクターの関数呼び出し演算子は、基本的に次のようになります。

__host__ __device__ bool operator()(float x) {
    bool mark = x < 0.0f;
    if (mark) {
        if (left neighbor of x > 1.0f) return false;
        if (right neighbor of x > 1.0f) return false;
        if (top neighbor of x > 1.0f) return false;
        //etc.
    }
    return mark;
}

現在、最初に CUDA カーネル (ネイバーに簡単にアクセスできる) を起動して、要素を適切にマークするという回避策を使用しています。その後、マークされた要素を Thrust の copy_if に渡し、マークされた要素のインデックスを抽出します。

処理された要素のインデックスを取得するために、threadIdx と blockIdx を直接使用する代わりに、counting_iterator を見つけました。以下の解決策を試しましたが、コンパイルすると、「/usr/include/cuda/thrust/detail/device/cuda/copy_if.inl(151): Error: Unaligned memory accesses not supported」というメッセージが表示されます。私の知る限り、整列されていない方法でメモリにアクセスしようとしているわけではありません。何が起こっているのか、および/またはこれを修正する方法を知っている人はいますか?

struct IsEmpty2 {
    float* xi;

    IsEmpty2(float* pXi) { xi = pXi; }

    __host__ __device__ bool operator()(thrust::tuple<float, int> t) {
        bool mark = thrust::get<0>(t) < -0.01f;
        if (mark) {
            int countindex = thrust::get<1>(t);
            if (xi[countindex] > 1.01f) return false;
            //etc.
        }
        return mark;
    }
};


thrust::copy_if(indices.begin(),
                indices.end(),
                thrust::make_zip_iterator(thrust::make_tuple(xi, thrust::counting_iterator<int>())),
                indicesEmptied.begin(),
                IsEmpty2(rawXi));
4

1 に答える 1

1

@phoad: 共有メモリについてはあなたの言うとおりです。既に返信を投稿した後、キャッシュがおそらく役立つだろうと考えて、私を驚かせました。しかし、あなたは迅速な対応で私を打ち負かしました。ただし、if ステートメントはすべてのケースの 5% 未満で実行されるため、共有メモリを使用するか、キャッシュに依存するかのいずれかが、パフォーマンスにほとんど影響を与えない可能性があります。

タプルは 10 個の値しかサポートしないため、3D の場合は 26 個の値に対してタプルのタプルが必要になります。タプルと zip_iterator の操作はすでに非常に面倒だったので、このオプションをパスします (コードの読みやすさの観点からも)。デバイス関数でthreadIdx.xなどを直接使用して提案を試みましたが、Thrustはそれを好みません。説明のつかない結果が得られているようで、Thrust エラーが発生することがあります。たとえば、次のプログラムは、最初に「Processing 10」から「Processing 41」を正しく出力しますが、「unspecified launch failure」を含む「thrust::system::system_error」を生成します。

struct printf_functor {
    __host__ __device__ void operator()(int e) {
        printf("Processing %d\n", threadIdx.x);
    }
};

int main() {
    thrust::device_vector<int> dVec(32);
    for (int i = 0; i < 32; ++i)
        dVec[i] = i + 10;

    thrust::for_each(dVec.begin(), dVec.end(), printf_functor());

    return 0;
}

blockIdx.x の印刷にも同じことが当てはまりますが、blockDim.x を印刷してもエラーは発生しません。私はクリーンな解決策を望んでいましたが、現在の回避策に行き詰まっていると思います。

于 2012-10-08T10:19:00.970 に答える