grep - count keyword in file - parallel program using CUDA resources.

countKeyword.cu

/*
	This program is to count keyword in file. The original code can run as command like:
		" ./grep keyword filename "
	in linux terminal. I modified it to let it run in parallel by using CUDA resources. 
	Now it can run in Visual Studio in machine with GPU device which support CUDA library.
	
 */

#include <iostream>
#include <sstream>
#include <fstream>
#include <bitset>
#include <string>
#include <unordered_map>
#include <chrono>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#ifndef __CUDACC__
#define __CUDACC__
#endif
#include <device_functions.h>

//#define MAX_DEF 4096
const int ntpb = 64 ; // number of threads per block

using namespace std::chrono;

using namespace std;

__global__ void grepx(char* input, char keychar, int* recordexist, int keywordidx) {
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	int t = threadIdx.x;

	//Global memory
	//if (input[i] == keychar) {
	//	recordexist[i] = keywordidx;
	//	
	//}

	//Shared memory
	__shared__ char s_input[ntpb];
	s_input[t] = input[i];
	__syncthreads();

	__shared__ int s_record[ntpb];
	s_record[t] = recordexist[i];
	__syncthreads();

	if (input[i] == keychar) {
		s_record[t] = keywordidx;
		__syncthreads();
	}

	recordexist[i] = s_record[t];
	__syncthreads();
}

__global__ void initBand(int* recordexist) {
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	int t = threadIdx.x;
	recordexist[i] = 0;
}

int grepPara(char* input, char* keywd, int inputLength, int keywordlen) {
	int countExist = 0;

	//block number
	int nblocks = (inputLength + ntpb - 1) / ntpb;
	int totalThread = nblocks*ntpb;

	//display block,thread info
	cout << "\nnblocks: " << nblocks << endl;
	cout << "ntpb: " << ntpb << endl;
	cout << endl;

	//creat variable for device, dinput_a is in and dinput_b is out for kernel
	char* dinput_a;
	int* dinput_b;
	char* dkeyword;
	string keyword = keywd;
	int* dkeyInt;

	//allocate memory
	cudaMalloc((void**)&dinput_a, inputLength * sizeof(char));
	cudaMalloc((void**)&dinput_b, inputLength * sizeof(int));
	cudaMalloc((void**)&dkeyword, keywordlen * sizeof(char));
	cudaMalloc((void**)&dkeyInt, keywordlen * sizeof(int));

	//copy from host to device
	cudaMemcpy(dinput_a, input, inputLength * sizeof(char), cudaMemcpyHostToDevice);
	cudaMemcpy(dkeyword, keywd, keywordlen * sizeof(char), cudaMemcpyHostToDevice);

	int* counter = new int();
	initBand << <nblocks, ntpb >> > (dinput_b);
	initBand << <1, keywordlen >> > (dkeyInt);

	for (int i = 0; i < keywordlen; i++) {
		char keychar = '\0';
		keychar = keywd[i];
		grepx << <1, keywordlen >> > (dkeyword, keychar, dkeyInt, (i + 1));
		grepx << <nblocks, ntpb >> > (dinput_a, keychar, dinput_b, (i + 1));
	}

	int* hinput_c = new int[inputLength];
	int* hkeyInt = new int[keywordlen];
	cudaMemcpy(hinput_c, dinput_b, inputLength * sizeof(int), cudaMemcpyDeviceToHost);
	cudaMemcpy(hkeyInt, dkeyInt, keywordlen * sizeof(int), cudaMemcpyDeviceToHost);

	for (int i = 0; i < (inputLength - keywordlen); i++) {
		bool match = true;
		for (int j = 0; j < keywordlen; j++) {
			if (hinput_c[i + j] != hkeyInt[j]) {
				match = false;
			}
		}
		if (match) {
			countExist++;
		}
	}

	cout << "Test!" << endl;

	//display
	//cout << endl;
	//cout << "hinput_c: ";
	//for (int i = 0; i < inputLength; i++){
	//	cout << hinput_c[i];
	//}
	//cout << endl;


	// free device memory
	cudaFree(dinput_a);
	cudaFree(dinput_b);
	cudaFree(dkeyword);

	cudaDeviceReset();	

	delete[] hinput_c;

	return countExist;
}


int grep(std::string input, std::string keyword, int inputLen, int keywordLen)
{
	int counter = 0;
	char* temp;
	temp = new char(keywordLen);

	for (int i = 0; i < (inputLen - keywordLen); i++) {
		for (int j = 0; j<keywordLen; j++) {
			temp[j] = input[i + j];
		}
		temp[keywordLen] = '\0';
		if (temp == keyword) {
			counter++;
		}
	}

	return counter;
}

static void show_usage()
{
	cerr << "Usage: \n"
		<< "This program is to count keyword in file.\n"
		<< "    grep keyword filename\n"
		<< endl;
}

int main(int argc, char* argv[]) {
	streampos size;

	char * memblock;

	if (argc < 2)
	{
		show_usage();
		return(1);
	}
	ifstream file(argv[2], ios::in | ios::binary | ios::ate);
	if (file.is_open())
	{
		size = file.tellg();
		memblock = new char[size];
		file.seekg(0, ios::beg);
		file.read(memblock, size);
		file.close();

		//string input = convert_char_to_string(memblock, size);
		string input = memblock;
		int inputLength = input.length();
		char* keywordInChar = argv[1];
		string keyword = argv[1];
		int keywordlen = keyword.length();

		cout << "grep - count keyword: '" << keyword << "' in file: '" << argv[2] << "'" << endl;
		cout << endl;
		//cout << "File input: " << input << endl;
		cout << "inputLength : " << inputLength << endl;
		cout << "keyword: " << keyword << endl;
		cout << "keywordlen: " << keywordlen << endl;
		cout << endl;

		//Serial grep function
		//int countSerial = grep(memblock, keywordInChar, inputLength, keywordlen);

		//Paraller grep function
		int countPara = grepPara(memblock, keywordInChar, inputLength, keywordlen);

		//cout << "word count(serial): " << countSerial << endl;
		cout << "word count(paraller): " << countPara << endl;

		cout << endl;
	}
	else {
		cout << "Unable to open file." << endl;
		show_usage();
	}



	//delete[] memblock;
	return 0;
}