CUB wählt aus, wenn mit zurückgegebenen Indexen.

Question

Ich habe in letzter Zeit Leistungsprobleme beim Verwenden der Thrust-Bibliothek festgestellt. Diese entstehen, wenn thrust Speicher im Basisteil einer großen verschachtelten Schleifenstruktur zuweist. Das ist offensichtlich unerwünscht, denn eine ideale Ausführung würde die Verwendung eines vorab allokierten Blocks globalen Speichers erfordern. Ich würde gerne den fehlerhaften Code auf eine der folgenden drei Arten entfernen oder verbessern:

1. Implementierung eines benutzerdefinierten thrust-Speicherallocators
2. Ersetzen des thrust-Codes durch CUB-Code (mit vorab allokiertem temporären Speicher)
3. Schreiben eines benutzerdefinierten Kernels, um das Gewünschte zu erreichen

Obwohl die dritte Option normalerweise meine bevorzugte Wahl wäre, handelt es sich um eine copy_if/select_if-Art von Operation, bei der sowohl die Daten als auch die Indizes zurückgegeben werden. Das Schreiben eines benutzerdefinierten Kernels wäre wahrscheinlich eine unnötige Wiederholung und daher würde ich eine der anderen beiden Optionen bevorzugen.

Ich habe viel Gutes über CUB gehört und sehe dies daher als ideale Gelegenheit, es wirklich einzusetzen. Was ich gerne wissen würde, ist:

Wie implementiert man ein CUB select_if mit zurückgegebenen Indizes?

Kann dies mit einem ArgIndexInputIterator und einem Funktor wie folgt erreicht werden?

struct GreaterThan
{
    int vergleich;

    __host__ __device__ __forceinline__
    GreaterThan(int vergleich) : vergleich(vergleich) {}

    __host__ __device__ __forceinline__
    bool operator()(const cub::ArgIndexInputIterator &a) const {
        return (a.value > vergleich);
    }
};

mit dem folgenden Code im Hauptteil des Codes:

//d_in = Geräteintarray
//d_temp_storage = ein vorab allokiierter Block

int schwellenwert;
GreaterThan select_op(schwellenwert);

cub::ArgIndexInputIterator input_itr(d_in);
cub::ArgIndexInputIterator output_itr(d_out); //????

CubDebugExit(DeviceSelect::If(d_temp_storage, temp_storage_bytes, input_itr, output_itr, d_num_selected, num_items, select_op));

Wird dies versuchen, unter der Haube irgendwelche Speicherzuweisungen vorzunehmen?

BEARBEITEN:

Also basierend auf Robert Crovellas Kommentar müsste der Funktor das Produkt aus dem Dereferenzieren eines cub::ArgIndexInputIterator nehmen, was ein cub::ItemOffsetPair sein sollte, wodurch der Funktor jetzt folgendermaßen aussieht:

struct GreaterThan
{
    int vergleich;

    __host__ __device__ __forceinline__
    GreaterThan(int vergleich) : vergleich(vergleich) {}

    __host__ __device__ __forceinline__
    bool operator()(const cub::ItemOffsetPair &a) const {
        return (a.value > vergleich);
    }
};

und im Code sollte d_out ein Gerätearray von cub::ItemOffsetPair sein:

//d_in = Geräteintarray
//d_temp_storage = ein vorab allokiierter Block

cub::ItemOffsetPair * d_out;
//d_out zuweisen

int schwellenwert;
GreaterThan select_op(schwellenwert);

cub::ArgIndexInputIterator input_itr(d_in);
CubDebugExit(DeviceSelect::If(d_temp_storage, temp_storage_bytes, input_itr, d_out, d_num_selected, num_items, select_op));

Answer 1

Nach etwas Herumprobieren und Nachfragen konnte ich einen einfachen Code entlang der von Ihnen vorgeschlagenen Linien zum Laufen bringen:

$ cat t348.cu
#include 
#include 
#define DSIZE 6

struct GreaterThan
{

    __host__ __device__ __forceinline__
    bool operator()(const cub::ItemOffsetPair &a) const {
        return (a.value > DSIZE/2);
    }
};

int main(){

  int num_items = DSIZE;
  int *d_in;
  cub::ItemOffsetPair * d_out;
  int *d_num_selected;
  int *d_temp_storage = NULL;
  size_t temp_storage_bytes = 0;

  cudaMalloc((void **)&d_in, num_items*sizeof(int));
  cudaMalloc((void **)&d_num_selected, sizeof(int));
  cudaMalloc((void **)&d_out, num_items*sizeof(cub::ItemOffsetPair));

  int h_in[DSIZE] = {5, 4, 3, 2, 1, 0};
  cudaMemcpy(d_in, h_in, num_items*sizeof(int), cudaMemcpyHostToDevice);

  cub::ArgIndexInputIterator input_itr(d_in);

  cub::DeviceSelect::If(d_temp_storage, temp_storage_bytes, input_itr, d_out, d_num_selected, num_items, GreaterThan());

  cudaMalloc(&d_temp_storage, temp_storage_bytes);

  cub::DeviceSelect::If(d_temp_storage, temp_storage_bytes, input_itr, d_out, d_num_selected, num_items, GreaterThan());
  int h_num_selected = 0;
  cudaMemcpy(&h_num_selected, d_num_selected, sizeof(int), cudaMemcpyDeviceToHost);
  cub::ItemOffsetPair h_out[h_num_selected];
  cudaMemcpy(h_out, d_out, h_num_selected*sizeof(cub::ItemOffsetPair), cudaMemcpyDeviceToHost);
  for (int i =0 ; i < h_num_selected; i++)
    printf("index: %d, offset: %d, value: %d\n", i, h_out[i].offset, h_out[i].value);

  return 0;
}
$ nvcc -arch=sm_20 -o t348 t348.cu
$ ./t348
index: 0, offset: 0, value: 5
index: 1, offset: 1, value: 4
$

RHEL 6.2, cub v1.2.2, CUDA 5.5

Answer 2

Ich bin in letzter Zeit auf Leistungsprobleme gestoßen, wenn ich die Thrust-Bibliothek verwende. Diese ergeben sich daraus, dass Thrust Speicher in der Basis einer großen verschachtelten Schleifenstruktur allokiert. Dies ist offensichtlich unerwünscht, da eine ideale Ausführung einen vorab allokierten Block von globalem Speicher verwendet.

Thrust ermöglicht es Ihnen, anzupassen, wie temporärer Speicher während der Algorithmusausführung allokiert wird.

Sehen Sie sich das custom_temporary_allocation Beispiel an, um zu sehen, wie Sie einen Cache für Ihren vorab allokierten Block erstellen können.