Changeset 5213 for icGREP/icgrep-devel


Ignore:
Timestamp:
Nov 10, 2016, 11:11:58 PM (3 years ago)
Author:
lindanl
Message:

Some fixes for editd. Add timer.

Location:
icGREP/icgrep-devel/icgrep/editd
Files:
2 edited

Legend:

Unmodified
Added
Removed
  • icGREP/icgrep-devel/icgrep/editd/EditdCudaDriver.h

    r5212 r5213  
    1616
    1717/// main - Program entry point
    18 ulong * RunPTX(std::string PTXFilename, char * fileBuffer, ulong filesize, const char * patternStr, unsigned patternLen) {
     18ulong * RunPTX(std::string PTXFilename, char * fileBuffer, ulong filesize, const char * patternStr, unsigned patternLen, int dist) {
    1919 
    2020  CUdevice    device;
     
    6868  int strides = filesize/(strideSize * 2) + 1;
    6969  int bufferSize = strides * strideSize;
    70   int outputSize = sizeof(ulong) * GROUPTHREADS * strides * 3 * GROUPBLOCKS;
     70  int outputSize = sizeof(ulong) * GROUPTHREADS * strides * (dist + 1) * GROUPBLOCKS;
    7171
    7272  checkCudaErrors(cuMemAlloc(&devBufferInput, bufferSize));
     
    9595  // std::cout << "Launching kernel\n";
    9696
     97  CUevent start;
     98  CUevent stop;
     99  float elapsedTime;
     100
     101  cuEventCreate(&start, CU_EVENT_BLOCKING_SYNC);
     102  cuEventRecord(start,0);
     103
    97104  // Kernel launch
    98105  checkCudaErrors(cuLaunchKernel(function, gridSizeX, gridSizeY, gridSizeZ,
    99106                                 blockSizeX, blockSizeY, blockSizeZ,
    100107                                 0, NULL, KernelParams, NULL));
    101   // std::cout << "kernel success.\n";
     108
     109  cuEventCreate(&stop, CU_EVENT_BLOCKING_SYNC);
     110  cuEventRecord(stop,0);
     111  cuEventSynchronize(stop);
     112
     113  cuEventElapsedTime(&elapsedTime, start, stop);
     114  printf("Elapsed time : %f ms\n" ,elapsedTime);
     115
    102116  // Retrieve device data
    103 
    104117  ulong * matchRslt = (ulong *) malloc(outputSize);
    105118  checkCudaErrors(cuMemcpyDtoH(matchRslt, devBufferOutput, outputSize));
  • icGREP/icgrep-devel/icgrep/editd/editd.cpp

    r5212 r5213  
    425425#define GROUPBLOCKS 64
    426426
    427 void editdGPUCodeGen(){ 
     427void editdGPUCodeGen(unsigned patternLen){ 
    428428    LLVMContext TheContext;
    429429    Module * M = new Module("editd-gpu", TheContext);
     
    445445    ExternalFileBuffer ResultStream(iBuilder, StreamSetType(iBuilder, editDistance+1, 1), addrSpace);
    446446
    447     const unsigned patternLen = 19;
    448447    kernel::editdGPUKernel editdk(iBuilder, editDistance, patternLen);
    449448    editdk.generateKernel({&CCStream}, {&ResultStream});
     
    512511    Type * const size_ty = iBuilder->getSizeTy();
    513512    Type * const voidTy = Type::getVoidTy(M->getContext());
    514     Type * const inputType = PointerType::get(ArrayType::get(mBitBlockType, 4), 0);
     513    Type * const inputType = PointerType::get(ArrayType::get(mBitBlockType, editDistance+1), 0);
    515514
    516515    ExternalFileBuffer MatchResults(iBuilder, StreamSetType(iBuilder, editDistance+1, 1));
     
    573572#ifdef CUDA_ENABLED 
    574573    setNVPTXOption();   
    575     if(codegen::NVPTX){   
    576         editdGPUCodeGen();
     574    if(codegen::NVPTX){
    577575
    578576        std::ifstream t(PatternFilename);
     
    581579            exit(-1);
    582580        } 
    583         std::string pattern_str((std::istreambuf_iterator<char>(t)), std::istreambuf_iterator<char>());
    584 
    585         ulong * rslt = RunPTX(PTXFilename, chStream, size, pattern_str.c_str(), pattern_str.length());
     581        std::string patterns((std::istreambuf_iterator<char>(t)), std::istreambuf_iterator<char>());
     582
     583        editdGPUCodeGen(patterns.length()/GROUPTHREADS - 1);
     584
     585        ulong * rslt = RunPTX(PTXFilename, chStream, size, patterns.c_str(), patterns.length(), editDistance);
    586586
    587587        editdFunctionType editd_ptr = editdScanCPUCodeGen();
Note: See TracChangeset for help on using the changeset viewer.