on-memory database の検索 : SELECT WHERE ～

週末の退屈しのぎにコード書き。

struct record {
  float height;
  float weight;
};

なんてな struct の配列 records[N] から weight < 0.1 を満たすものを列挙します。
SQLでいうところの

SELECT * 
  FROM records
  WHERE weight < 0.1;

ですわな。
これをCUDAでやるとなると records[N] および 検索結果を格納するindex列 int indices[N] はDevice上に置くことになります。
WHERE 節に与える条件式は　device-lambda 使って楽しました。

// nvcc fake_DB.cpp --expt-extended-lambda 
// CUDA 8.0

#include <cuda_runtime.h>
#include <device_launch_parameters.h>

struct record {
  float height;
  float weight;
};

template<typename Where>
__global__ void kernel_select(
    const record* records, unsigned int size,
    Where condition, int* count, int* indices) {

  unsigned int i = blockDim.x * blockIdx.x + threadIdx.x;
  if (i == 0) *count = 0; // 一回だけ、 *count を 0-clear
  __syncthreads();

  while (i < size) {
    // i番目のrecordが　WHERE条件を満たしているなら
    // indicesにiを追加
    if ( condition(records[i]) ) {
      indices[atomicAdd(count, 1)] = i;
    }
    i += gridDim.x * blockDim.x;
  }
}

#include <random>
#include <iostream>
#include <iomanip>
#include <algorithm>
using namespace std;

int main() {
  const int N = 100;
  record records[N];

  // fill records[] with random numbers
  mt19937 gen;
  uniform_real_distribution<float> dist;
  generate_n(records, N, [&]() { return record{ dist(gen),dist(gen) }; });

  // result
  int indices[N];
  int count;

  // allocate device-mem.
  record* dev_records;
  cudaMalloc(&dev_records, N * sizeof(record));
  int* dev_indices;
  cudaMalloc(&dev_indices, N * sizeof(int));
  int* dev_count;
  cudaMalloc(&dev_count, sizeof(int));

  // copy records from HOST to DEVICE
  cudaMemcpy(dev_records, records, N * sizeof(record), cudaMemcpyHostToDevice);

  // SELECT * FROM dev_records
  // WHERE weight < 0.1
  auto where = [] __device__(const record& rec) { return rec.weight < 0.1f; };

  dim3 grid = (N + 255) / 256;
  dim3 block = 256;
  kernel_select <<<grid, block >> >(
      dev_records, N,
      where,
      dev_count, dev_indices);

  // copy result from DEVICE to HOST.
  cudaMemcpy(&count, dev_count, sizeof(int), cudaMemcpyDeviceToHost);
  cudaMemcpy(indices, dev_indices, count * sizeof(int), cudaMemcpyDeviceToHost);
//*/
  for (int i = 0; i < count; ++i) {
    cout << left 
         << setw(10) << records[indices[i]].height << " : "
         << setw(10) << records[indices[i]].weight << endl;
  }
//*/
  cout << count << " records found." << endl;

  cudaFree(dev_records);
  cudaFree(dev_count);
  cudaFree(dev_indices);
  cudaDeviceReset();

}

メモリの乏しいDeviceでは扱えるrecord数をそんなに大きくはできません。まぢに使うとなればrecords[N]を分割し、Deviceに乗せては検索を繰り返すことになるでしょうね。