Skip to content

Instantly share code, notes, and snippets.

Embed
What would you like to do?
grep - count keyword in file - parallel program using CUDA resources.
/*
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;
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
You can’t perform that action at this time.