CUDAを使ったプログラミング データ検索

#include <stdio.h>
#include <iostream>
#include <time.h>

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "cuda01.cuh"

// GPUを使った検索(カーネル)
__global__ void gpu_search(int* values, int* results)
{
	int idx = blockIdx.x * blockDim.x + threadIdx.x;	// インデックスを計算
	if (values[idx] % 5 == 0)
		results[idx] = values[idx];
}

// CPUを使った検索
void cpu_search(int n, int* values, int* results)
{
	for (int i = 0; i < n; i++) {
		if (values[i] % 5 == 0) {
			results[i] = values[i];
		}
	}
}

// メイン処理(外部cppから呼び出す)
int cuda_call(void)
{
	int N = 256 * 4; //256の倍数にしないと正常に動作しない
	int* host_values, * host_results, * dev_values, * dev_results;

	// CPU側の領域確保
	host_values  = (int*)malloc(N * sizeof(int));
	host_results = (int*)malloc(N * sizeof(int));

	// 連番生成
	for (int i = 0; i < N; i++) {
		host_values[i] = i;
	}
	memset(host_results, 0, N * sizeof(int));	// 戻り値配列を0初期化

	// CPUで計算
	int start_cpu = clock();	// CPUタイマー
	cpu_search(N, host_values, host_results);
	int ms_cpu = clock() - start_cpu;

	// GPUで計算
	int start_gpu = clock();	// GPUタイマー
	// デバイス(GPU)側の領域確保
	cudaMalloc(&dev_values,  N * sizeof(int));
	cudaMalloc(&dev_results, N * sizeof(int));

	// CPU ⇒ GPUのデータコピー
	cudaMemcpy(dev_values,  host_values,  N * sizeof(int), cudaMemcpyHostToDevice);
	cudaMemcpy(dev_results, host_results, N * sizeof(int), cudaMemcpyHostToDevice);

	// GPUで計算
	gpu_search <<< (N + 255) / 256, 256 >>> (dev_values, dev_results);

	// GPU ⇒ CPUのデータコピー
	cudaMemcpy(host_results, dev_results, N * sizeof(int), cudaMemcpyDeviceToHost);
	int ms_gpu = clock() - start_cpu;

	// 戻り値を出力
	std::cout << "result:" << std::endl;
	for (int j = 0; j < N; j++) {
		std::cout << host_results[j] << ",";
	}
	std::cout << std::endl;

	// 計測時間を出力
	std::cout << "CPU: " << ms_cpu << "[ms]" << ", GPU: " << ms_gpu << "[ms]" << std::endl;
	return 0;
}

__global__ void gpu_search(int* values, int* results)
{
	int idx = blockIdx.x * blockDim.x + threadIdx.x;	// インデックスを計算
	int ofs = 0;
	if (values[idx] % 5 == 0) {
		while (true) {
			if (results[ofs] == 0) {
				results[ofs] = values[idx];
				break;
			}
			else {
				ofs++;
				if (ofs >= 1024)
					break;
			}
			__syncthreads();
		}
	}
}

// GPUを使った検索(カーネル)
__global__ void gpu_search(int* values, int* results, int* flags)
{
	int idx = blockIdx.x * blockDim.x + threadIdx.x;	// インデックスを計算
	int ofs = 0;
	while (flags[idx] != 1) {
		__syncthreads();
	}
	if (values[idx] % 5 == 0) {
		while (true) {
			if (results[ofs] == 0) {
				results[ofs] = values[idx];
				break;
			}
			else {
				ofs++;
				if (ofs >= 1024)
					break;
			}
		}
	}
	flags[idx + 1] = 1;
}

#define BUF_SIZE (4000)
#define BATCH_SIZE (256)

// GPUを使った検索(カーネル)
__global__ void gpu_search(int* values, int* results, int* flags, int *pos)
{
	int idx = blockIdx.x * blockDim.x + threadIdx.x;	// インデックスを計算
	int ofs = 0;
	int i = 0;
	if (idx % BATCH_SIZE != 0)
		return;

	for (i = 0; i < BATCH_SIZE; i++) {
		if (values[idx + i] % 5 == 0) {
			results[idx + ofs] = values[idx + i];
			ofs++;
		}
	}

	if (idx == 0) {
		pos[0] = ofs;
		flags[idx + BATCH_SIZE] = 1;
		return;
	}
	while (flags[idx] != 1) {
		__syncthreads();
	}
	for (i = 0; i < BATCH_SIZE; i++) {
		if (results[idx + i] != 0) {
			results[pos[0]] = results[idx + i];
			pos[0]++;
		}
		else {
			break;
		}
	}

	flags[idx + BATCH_SIZE] = 1;	
}