den ersten Index eines jeden eindeutigen Wert in CUDA-Array zu finden

Question

Angenommen wir ein Array wie dieses:

0, 0, 0, 1, 2, 2, 2, 3, 3, 4, ... 

Ich möchte den Index jedes erste Vorkommen eines jeden Wert haben, also in diesem Beispiel [0 , 3, 4, 7, 9]. Das Array ist sortiert und alle möglichen Werte sind bekannt und fortlaufend.

Mögliche Lösungen, die ich habe verwendet einen Kernel für jedes Element in diesem Array und verwenden Sie eine Atommin, um den niedrigsten Index zu speichern. Aber ich nehme an, dass ein besserer Ansatz möglich ist.

Answer 1

Sie können dies tun, mit ein einzelner Aufruf an thrust::unique_by_key(), wenn Sie einen Vektor von Indizes z über thrust::sequence(). Hier ist ein Beispiel gearbeitet:

$ cat t3.cu 
#include <thrust/device_vector.h> 
#include <thrust/copy.h> 
#include <thrust/unique.h> 
#include <thrust/sequence.h> 
#include <iostream> 

int main(){ 

    int keys[] = {0, 0, 0, 1, 2, 2, 2, 3, 3, 4}; 
    int ks = sizeof(keys)/sizeof(keys[0]); 
    thrust::device_vector<int> d_keys(keys, keys+ks); 
    thrust::device_vector<int> d_result(ks); 
    thrust::sequence(d_result.begin(), d_result.end()); 
    int rs = (thrust::unique_by_key(d_keys.begin(), d_keys.end(), d_result.begin())).first - d_keys.begin(); 
    thrust::copy_n(d_result.begin(), rs, std::ostream_iterator<int>(std::cout, ",")); 
    std::cout << std::endl; 
} 

$ nvcc -arch=sm_35 -o t3 t3.cu 
$ ./t3 
0,3,4,7,9, 
$ 

Die wichtige Aktivität hier auftritt, ist Strom Verdichtung und Schub liefert eine nice set of routines für verschiedene Anwendungsfälle. Zum Beispiel könnte diese Operation auch mit thrust::unique_copy() durchgeführt werden, und in diesem Fall könnten Sie mit etwas zusätzlicher Code-Komplexität die Notwendigkeit für den thrust::sequence() Aufruf (es würde durch eine thrust::counting_iterator gezippt zusammen mit Ihren Daten und einem entsprechenden Auswahlfunktor ersetzt werden) , benötigt aber immer noch einen Ausgangsvektor gleicher Länge.

Answer 2

Wie @tera darauf hingewiesen hat, können Sie eine Zahl mit der vorherigen Zahl vergleichen, um festzustellen, ob es sich um das erste Vorkommen in einer Folge eindeutiger Zahlen handelt. Sie können einen Kernel schreiben, um eine Maske für dieses Kriterium zu erzeugen, so dass das Masken-Array den Index für eine Zahl enthält, die ein erstes Vorkommen und eine negative Zahl ist (wie -1, da es kein Index sein kann). Verwenden Sie danach Schub, um die Nicht-1-Werte zu zählen, indem Sie ein Prädikat verwenden. Kopieren Sie dann diese Werte aus der Maske mit dem gleichen Prädikat wie oben. Abschließend kopieren Sie die Ergebnisse zurück zum Host.

Hier ist eine Beispielimplementierung des oben genannten Ansatzes.

#include <iostream> 
#include <cuda_runtime.h> 
#include <thrust/device_vector.h> 
#include <thrust/count.h> 
#include <thrust/copy.h> 

using namespace std; 


//Copy index 
__global__ void is_first_occurence(int* input, int* is, int count) 
{ 
    const int tid = blockIdx.x * blockDim.x + threadIdx.x; 

    if(tid<count) 
    { 
     if(tid == 0) 
     { 
      is[0] = 0; 
     } 
     else if(input[tid] != input[tid-1]) 
     { 
      is[tid] = tid; 
     } 
     else 
     { 
      is[tid] = -1; 
     } 
    } 
} 


struct isFirst 
{ 
    __host__ __device__ bool operator()(const int x) 
    { 
    return (x != -1); 
    } 
}; 

int main(int argc, char** argv) 
{ 
    const int count = 13; 

    std::vector<int> arr = { 0, 0, 0, 1, 1, 2, 2, 2, 3, 3, 4, 4 ,4 }; 

    thrust::device_vector<int> arr_d = arr; 
    thrust::device_vector<int> mask_d(arr_d.size()); 

    int* pArr = thrust::raw_pointer_cast(arr_d.data()); 
    int* pMask = thrust::raw_pointer_cast(mask_d.data()); 

    dim3 block(16); 
    dim3 grid((count + block.x -1)/block.x); 

    is_first_occurence<<<grid,block>>>(pArr, pMask, count); 
    cudaDeviceSynchronize(); 

    int num_unique = thrust::count_if(mask_d.begin(), mask_d.end(), isFirst()); 

    thrust::copy_if(mask_d.begin(), mask_d.end(), arr_d.begin(), isFirst()); 

    std::vector<int> unique_indices(num_unique); 

    thrust::copy(arr_d.begin(), arr_d.begin() + num_unique, unique_indices.begin()); 

    for(auto i:unique_indices) 
    { 
     cout<<i<<endl; 
    } 

    return 0; 
} 

kompiliert und getestet mit dem folgenden Befehl:

nvcc -o get_unique get_unique.cu -std = C++ 11 -arch = sm_61