快速 select 算法在 CUDA 上处理运行的只读输入

Question

我正在寻找一种类似于 quick select 的算法，它不会改变其输入。我更喜欢只读数据，这样在 CUDA GPU 上执行时内存效率更高。

Answer 1

您可能想避免在 GPU 上分配临时内存，大多数合理排序 and/or select 算法会根据您的语句 "so it will be more memory efficient to execute on CUDA GPU".

首先让我说我认为这可能是个坏主意，因为尽管您可能已经满足了 "memory efficiency" 的想法，但这样做几乎肯定会削弱任何算法的性能。给排序或 select 算法临时 space 来跟踪它所做的工作可能会产生更快的代码，但会以上述临时内存分配为代价。

有了这个序言，这里有一种可能的方法：一种强力算法，它对只读数组中的元素进行排序，直到找到满足给定 selected 项的元素。基本思想非常简单：对于只读输入数组中的每个项目，计算 entire 数组中位于它之上的项目数（对于降序排序）。重复项需要特殊情况处理，但我已将其包括在内。找到 selected 项目后，可以取消搜索。

这是一个完整的示例，带有结果验证：

$ cat t729.cu
#include <stdio.h>
#include <stdlib.h>

// for validation
#include <algorithm>

#define DSIZE 10000
#define RG    1000
#define SELECT_ITEM 100

#define nBLK  64
#define nTPB  256

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)



#define DESCEND true
#define ASCEND  false



template <bool descending, typename T>
__global__ void my_select(const T *data, const unsigned length, const unsigned select, volatile int *index){

  int idx = threadIdx.x+blockDim.x*blockIdx.x;
  while ((*index == -1) && (idx < length)){
    T my_val = data[idx];
    unsigned my_index = 0;
    // count the number of values higher (or lower) than me
    for (unsigned i = 0; i < length; i++){
        T temp = data[i];
        if (descending){ // request to select item in descending order
          if (temp > my_val) my_index++;
          else if ((temp == my_val) && (i < idx)) my_index++; // handle duplicates
          }
        else {  // request to select item in ascending order
          if (temp < my_val) my_index++;
          else if ((temp == my_val) && (i < idx)) my_index++; // handle duplicates
          }}
    if (my_index == select) *index = idx;
    idx += blockDim.x*gridDim.x;
    }
}

int main(){

  int *h_data, *d_data, *d_result, h_result = -1;
  h_data = (int *)malloc(DSIZE*sizeof(int));
  if (h_data == NULL) {printf("malloc fail\n"); return -1;}
  cudaMalloc(&d_data, DSIZE*sizeof(int));
  cudaMalloc(&d_result, DSIZE*sizeof(int));
  cudaCheckErrors("cudaMalloc fail");
  for (int i =0; i < DSIZE; i++) h_data[i] = rand()%RG;
  cudaMemcpy(d_data, h_data, DSIZE*sizeof(int), cudaMemcpyHostToDevice);
  cudaMemcpy(d_result, &h_result, sizeof(int),   cudaMemcpyHostToDevice);
  cudaCheckErrors("cudaMemcpy fail");
  my_select<DESCEND><<<nBLK, nTPB>>>(d_data, DSIZE, SELECT_ITEM, d_result);
  cudaMemcpy(&h_result, d_result, sizeof(int), cudaMemcpyDeviceToHost);
  cudaCheckErrors("cudaMemcpy/kernel fail");
  printf("the item index at select position %d is %d, item value = %d\n", SELECT_ITEM, h_result, h_data[h_result]);
  // validation
  std::sort(h_data, h_data+DSIZE);
  printf("the item value at that sorted position is                %d\n", h_data[(DSIZE-1) - SELECT_ITEM]); // assumes DESCEND

  return 0;
}
$ nvcc -o t729 t729.cu
$ ./t729
the item index at select position 100 is 2858, item value = 990
the item value at that sorted position is                990
$

在我看来，更好的方法是复制您的输入数据，并将该副本传递给您选择的 GPU 优化 selection 算法，例如 this one 和让该算法使用它需要的任何临时 space 。这肯定会比上述方法更快，并且可能比您可能想出的不使用临时 space.

的任何方法都要快

快速 select 算法在 CUDA 上处理运行的只读输入

Quick select algorithm working on read only input to run on CUDA

sorting

algorithm

cuda

gpu

快速 select 算法在 CUDA 上处理 运行 的只读输入

Quick select algorithm working on read only input to run on CUDA

sorting

algorithm

cuda

gpu

快速 select 算法在 CUDA 上处理运行的只读输入