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

#define BLOCKDIM 1024
/*
#include <thrust/device_vector.h>
#include <thrust/reduce.h>
#include <thrust/functional.h>
#include <thrust/inner_product.h>
using namespace thrust;
using namespace thrust::placeholders;
*/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char *file, int line, bool abort=true)
{
    if (code != cudaSuccess)
    {
        fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        if (abort) exit(code);
    }
}

/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
///////////////////////////////Important Configuration/////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////

const char FileName[] = "01.nt"; // the file must have
const char OutputFileName[] = "v.txt"; // the file must have, but must be empty
const char QuestionFileName[] = "line1"; // the file must have


// Threadchucksize mean "Size" in byte of part of the file that will be sent to the gpu memory
// not worry about setting size higher than real file, because the program will auto adjust it
// Example , if your gpu memory limit at 500 MB but your file is large such as 5 GB , you should set threadchucksize to 500 * 1,000,000byte = 500,000,000
// and your part of file will be sent 5000 MB/500 MB = 10 times
// if you sent and found some memory error, it can be the os or other software use vram too, so decrease threadchucksize until no error, such as from 500,000,000 change to 200,000,000
// suggestion, you should set size as beautiful ten-end number as 100000, 20000000, 50000000
const long threadchucksize = 300000000;


// this Blocksize not mean thread per block but mean "size" of data chuck that each thread will compute from the big global data chuck (that locate in gpu)
// warning that allThreadInUse multiply with blocksize must higher than threadchucksize, if not it will incorrect result
// Example : if threadchucksize = 100,000,000 (aka 100MB. chuck of file sent to gpu) and allThreadInUse = 4096 and blocksize = 25000
// then you must check that 25000 * 4096 > 100000000 ? which is 102,400,000 > 100,000,000 so it true and can be use
// suggestion, you should set size as threadchucksize % blocksize = 0, it will be bug free.
const long blocksize = 10000;


const int NumberOfComputeBlock = 30; // aka gridsize
const int NumberOfThreadsPerBlock = 1024; //(rely on your gpu spec)

// allThreadInUse is all concurrent thread that run in the gpu,
// it can be higher than physical cuda core on gpu, because gpu can queue it and make you feel like it concurrent
// but if allThreadInUse is much higher, the answer vector that collect answer from each thread will be larger. so threadoff
const int allThreadInUse = NumberOfThreadsPerBlock * NumberOfComputeBlock;
const long sizeofAnswerVector = allThreadInUse;


const int sizeofQuestionArray = 2048; // maximum list size of question, such as = 2048 mean this program support maximum 2048 question
const int sizeofMaximumQuestionWord = 2048; // maximum string size of each question, such as = 2048 mean each question can't be larger than 2048 byte


/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////



int WordCount;
__device__ size_t d_strlen (const char *str)
{
    return (*str) ? d_strlen(++str) + 1 : 0;
}
__device__ int d_strncmp(const char *ptr0, const char *ptr1, size_t len)
{
    while(len--)
        if(*ptr0++!=*ptr1++)
            return *(unsigned char*)(ptr0 - 1) - *(unsigned char*)(ptr1 - 1);
    return 0;
}
__device__ unsigned int string_search(long start, long end, char* target, char *buffer) {
    unsigned int i;
    unsigned int found=0;
    //printf("Receiveing : target : %s  buffer : %s\n", target, buffer);
    //printf("My Buffer: ");
    //for (i=start;i <= end; i++)
    //    printf("%c", *(buffer+i));
    //printf("\n");
    for (i=start;i <= end; i++) {
        int t = d_strncmp(&buffer[i], target, d_strlen(target));
        //printf("t dkmfdsfdfdspfdsfpodsfjkdpsof: %d \n", t);
        if (t == 0) {
            //if (i <= overflowRegion)
            found++;
        }
    }
    //printf("overflowRegion : %d\n", overflowRegion);
    //printf("Receiveing Found : %d\n", found);
    return found;
}
__device__ unsigned int string_search_rr(long start, long end, char* target, char *buffer,int overflowStringSize, char options, char *changebuffer) {

    unsigned int i;
    unsigned int found=0;

    for (i=start;i <= end  ; i++) {
        int t = d_strncmp(&buffer[i], target, d_strlen(target));
        if (t == 0 ) {
            //if (i <= overflowRegion)
            found++;

            for (int j = i; j < i + d_strlen(target); ++j) {
                //printf("change at j : %d i : %d\n", j,i);
                *(changebuffer+j) = '$';
            }

        } else {
	    // if you want to search word in any file (not only file that contain one line per word. please comment this block)
	    //
	    int j = 0;
 	    for (j = i; j <= end; j++) {
	    	if (*(buffer+j) == '\n')
	 	    break;
	    }
	    i = j;
	    //
	    //-----------------------------------------------------------------------------------------------------------------
	}

    }

    return found;
}
__global__ void cuda_stringsearch (long bufferstart, long bufferend, char* target, char* buffer, int* allcount, int overflowStringSize, long *answerVector, char* changebuffer) {
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    //printf("hello from thread %d\n", index);
/*  
  if (index == 1) {
        printf("GPU KERNEL :: Hello from threads %d  Given word %s  size %d \n", index, target, d_strlen(target));
    }
*/
    //printf("Hello from threads %d  Given word %s\n", index, target);

   // long blocksize = 500/*50000*/;
    long extendblocksize = blocksize + overflowStringSize - 2;
    long startpoint = index * blocksize;
    long endpoint = startpoint + blocksize  - 1;

    if (startpoint <= bufferend) {
        if (endpoint > bufferend)
            endpoint = bufferend;

        //int count = 10;
        int count = string_search_rr(startpoint, endpoint,target, buffer, overflowStringSize, 'd', changebuffer);
        //printf("threads %d count %d  getting data :  startpoint %ld  endpoint %ld  overflowStringSize %d\n", index,count,startpoint, endpoint, overflowStringSize);
        *(answerVector + index) = count;
        //*allcount += count;
    }

}

long* createVector (long size, long inivalue) {
    long* vector = (long*) malloc(sizeof(long)*size);
    for (long i = 0; i < size; ++i) {
        vector[i] = inivalue;
    }
    return vector;
}
void readVector (long* vector, long size) {
    for (long i = 0; i < size; ++i) {
        printf("%ld ",vector[i]);
    }
    printf("\n");
}
long sumVector (long* vector, long size) {
    long sum = 0;
    for (long i = 0; i < size; ++i) {
        sum+= vector[i];
    }
    return sum;
}
int main(int argc, char **argv) {

    time_t timestart = time(NULL);

    FILE * pFile;
    long lSize;
    pFile = fopen ( FileName , "r" );
    if (pFile==NULL) {fputs ("File error",stderr); exit (1);}
    fseek (pFile , 0 , SEEK_END);
    lSize = ftell (pFile);
    rewind (pFile);
    long BufferSize = sizeof(char)*lSize;
    printf("Buffer index size %lu \n",BufferSize);

    int reverseoffset = 0;
    /* create threads */
    long endpoint = 0,startpoint = 0;




    FILE * questionFile;
    long lSizeQ;
    questionFile = fopen(QuestionFileName, "r");
    if (questionFile == NULL) {fputs ("File error", stderr); exit(1);}
    fseek(questionFile, 0, SEEK_END);
    lSizeQ = ftell(questionFile);
    rewind(questionFile);
    long QuestionBufferSize = sizeof(char)*lSizeQ;
    printf("Question Buffer index size %lu \n", QuestionBufferSize);
    char *Question_Buffer = (char*) malloc (lSizeQ);
    fread(Question_Buffer, 1, QuestionBufferSize, questionFile);
    printf("This is question file --------\n");

    //printf("%s\n", Question_Buffer);

    long start = 0, end = 0;
    int Question_maxLength = 0;
    char** questionArray = (char**) malloc(sizeof(char*)*sizeofQuestionArray);
    long* questionAnswer = (long*) malloc(sizeof(long)*sizeofQuestionArray); // found word list

    int questionCount = 0;
    for (int j = 0; j <= strlen(Question_Buffer); ++j) {
        end++;
        if (*(Question_Buffer+j) == '\n' || *(Question_Buffer+j) == '\0') {
            questionCount++;
            //piece = (char*) malloc(sizeof(char)*2048);
            //memcpy(piece, (Question_Buffer+start), end - start - 1);
            *(questionArray+questionCount) = (char*) malloc(sizeof(char)*sizeofAnswerVector);
            memcpy(*(questionArray+questionCount), (Question_Buffer+start), end - start - 1);
            *(questionAnswer+questionCount) = 0; /* each question start founded = zero */               /*(long) strlen(*(questionArray+questionCount)) this commented code use to check if for loop work!*/
            if (strlen(*(questionArray+questionCount)) > Question_maxLength)
                Question_maxLength = strlen(*(questionArray+questionCount));
            //printf("print piece %s|||\n", *(questionArray+questionCount));
            //printf("piece length : %lu \n", strlen(*(questionArray+questionCount)));
            start = end;

        }
    }
/*
    for (int k = 1; k <= questionCount; ++k) {
        printf("element at : %d is : %s value is : %lu  length is %zu\n",  k, *(questionArray+k), *(questionAnswer+k), strlen(*(questionArray+k)));
    }
*/
    printf("Question max length : %d\n", Question_maxLength);
    printf("Question elements count : %d\n", questionCount);
    free(Question_Buffer);
    printf("This is question file --------\n");

    int overflowStringSize = Question_maxLength /*- 1*/;
    printf("Overflow String size : %d\n", overflowStringSize);



    FILE * outputFile;
    long lSize2;
    outputFile = fopen(OutputFileName, "a");
    if (outputFile==NULL) {fputs ("File error",stderr); exit (1);}


    int count = 0;
    int* countPTR = &count;
    int overflowRegion = threadchucksize - 1;
    while (1){
        //printf("precount  all count %d\n", count);

        char *buffer;
        startpoint = 0;
        endpoint = threadchucksize + overflowStringSize - 1;
        buffer = (char*) malloc (sizeof(char)*(threadchucksize + overflowStringSize));
        fseek (pFile , reverseoffset , SEEK_CUR);
        reverseoffset = -1 * (overflowStringSize - 1);

        fread (buffer,1,endpoint,pFile);
        if (BufferSize <= threadchucksize)
            endpoint = BufferSize;
        printf("This will send buffer start at %ld to %ld of all %ld\n", startpoint, endpoint, BufferSize);

        //int j = 0;
        //count += string_search(startpoint, endpoint, DefineWord, buffer);
        //printf("%s||| count : %d\n", buffer, count);
        //printf("-------\n");




        char *dev_buffer;
        char *dev_changebuffer;
        int *dev_countPTR;
        //char *dev_defineword;
        //long *dev_answerVector;

        cudaMalloc((void**)&dev_buffer, sizeof(char)*(threadchucksize + overflowStringSize));
        cudaMalloc((void**)&dev_changebuffer, sizeof(char)*(threadchucksize + overflowStringSize));
        cudaMalloc((void**)&dev_countPTR, sizeof(int));

        cudaMemcpy(dev_buffer, buffer, sizeof(char)*(threadchucksize + overflowStringSize), cudaMemcpyHostToDevice);
        cudaMemcpy(dev_changebuffer, buffer, sizeof(char)*(threadchucksize + overflowStringSize), cudaMemcpyHostToDevice);

        cudaMemcpy(dev_countPTR, countPTR, sizeof(int),cudaMemcpyHostToDevice);

        for (int question = 1; question <= questionCount - 1; ++question) {
        //int question = 1;
            long size_answerVector = sizeofAnswerVector;
            long* answerVector = createVector(size_answerVector,0);
	    char *dev_defineword;
	    long *dev_answerVector;

//	    printf("HOST :: starting iteration %d at question : %s  string length : %zu\n",question, *(questionArray+question),  strlen(*(questionArray+question)));
/*		for (int d = 0; d < strlen(*(questionArray+question)); d++) {
			printf("%c",*(*(questionArray+question)+d));
		}
		printf("\nend test \n");*/
            cudaMalloc((void**)&dev_answerVector, sizeof(long)*size_answerVector);
            cudaMalloc((void**)&dev_defineword, /*sizeof(char)**/ sizeofMaximumQuestionWord/*strlen(*(questionArray+question))*/);

            cudaMemcpy(dev_answerVector, answerVector, sizeof(long)*size_answerVector, cudaMemcpyHostToDevice);
            cudaMemcpy(dev_defineword, *(questionArray+question), /*sizeof(char)**/ sizeofMaximumQuestionWord /*strlen(*(questionArray+question))*/, cudaMemcpyHostToDevice);
            //printf("iteration at question : %s\n", *(questionArray+question));

            cuda_stringsearch<<<NumberOfComputeBlock,NumberOfThreadsPerBlock>>>(startpoint, endpoint, dev_defineword, dev_buffer, dev_countPTR, overflowStringSize, dev_answerVector, dev_changebuffer);
            cudaDeviceSynchronize();
            cudaMemcpy (answerVector, dev_answerVector, sizeof(long)*size_answerVector, cudaMemcpyDeviceToHost);

            cudaFree(dev_answerVector);
            cudaFree(dev_defineword);
	    
            //readVector(answerVector, size_answerVector); //uncomment this to diagnostic answer vector matrix
            long iterationsum =  sumVector(answerVector, size_answerVector);

            *(questionAnswer+question) += iterationsum;

            printf("HOST :: Finish iteration %d at question : %s  temporary founded %ld\n\n",question, *(questionArray+question),  iterationsum);
            free(answerVector);
        }


        gpuErrchk( cudaPeekAtLastError() );
        gpuErrchk( cudaDeviceSynchronize() );

        cudaMemcpy (buffer, dev_changebuffer,sizeof(char)*(threadchucksize + overflowStringSize),cudaMemcpyDeviceToHost);
        //cudaMemcpy (countPTR, dev_countPTR, sizeof(int), cudaMemcpyDeviceToHost);
        cudaFree(dev_buffer); cudaFree(dev_countPTR); cudaFree(dev_changebuffer);

        printf("---------saving change buffered----------------------------------------------------------------\n");
        endpoint = threadchucksize - 1;
        if (endpoint > BufferSize)
            endpoint = BufferSize;
        //printf("startpoint : %ld     endpoint : %ld   BufferSize : %ld \n", startpoint, endpoint, BufferSize);
        for (int i=startpoint;i <= endpoint; i++) {
            if (*(buffer + i) != '\0' && *(buffer + i) != '$' )
                fprintf(outputFile, "%c", *(buffer + i));
            //printf("%c", *(buffer + i));
        }
        printf("\n");

        printf("-------------------------------------------------------------------------------------------\n");

        //fprintf(outputFile, "%s",buffer);

        BufferSize = BufferSize - threadchucksize;

        free(buffer);
        if (BufferSize <= 0)
            break;
    }
    for (int k = 1; k <= questionCount; ++k) {
        printf("element at : %d is : %s finally founded : %lu\n",  k, *(questionArray+k), *(questionAnswer+k));
    }
    fclose (pFile);
    fclose (outputFile);

    printf("\nestimate using time : %.2f\n", (double)(time(NULL) - timestart));

    return EXIT_SUCCESS;
}

	.file	"tmpxft_002e1359_00000000-6_cuda3.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB3646:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE3646:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.rodata._Z9gpuAssert9cudaErrorPKcib.str1.1,"aMS",@progbits,1
.LC0:
	.string	"GPUassert: %s %s %d\n"
	.section	.text._Z9gpuAssert9cudaErrorPKcib,"axG",@progbits,_Z9gpuAssert9cudaErrorPKcib,comdat
	.weak	_Z9gpuAssert9cudaErrorPKcib
	.type	_Z9gpuAssert9cudaErrorPKcib, @function
_Z9gpuAssert9cudaErrorPKcib:
.LFB3635:
	.cfi_startproc
	endbr64
	testl	%edi, %edi
	je	.L8
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	%ecx, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	%edi, %ebx
	subq	$24, %rsp
	.cfi_def_cfa_offset 48
	movl	%edx, 12(%rsp)
	movq	%rsi, (%rsp)
	call	cudaGetErrorString@PLT
	movl	12(%rsp), %r9d
	movq	(%rsp), %r8
	leaq	.LC0(%rip), %rdx
	movq	stderr(%rip), %rdi
	movq	%rax, %rcx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
	testb	%bpl, %bpl
	je	.L2
	movl	%ebx, %edi
	call	exit@PLT
.L2:
	addq	$24, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
.L8:
	.cfi_restore 3
	.cfi_restore 6
	ret
	.cfi_endproc
.LFE3635:
	.size	_Z9gpuAssert9cudaErrorPKcib, .-_Z9gpuAssert9cudaErrorPKcib
	.text
	.globl	_Z8d_strlenPKc
	.type	_Z8d_strlenPKc, @function
_Z8d_strlenPKc:
.LFB3636:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE3636:
	.size	_Z8d_strlenPKc, .-_Z8d_strlenPKc
	.globl	_Z9d_strncmpPKcS0_m
	.type	_Z9d_strncmpPKcS0_m, @function
_Z9d_strncmpPKcS0_m:
.LFB3637:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE3637:
	.size	_Z9d_strncmpPKcS0_m, .-_Z9d_strncmpPKcS0_m
	.globl	_Z13string_searchllPcS_
	.type	_Z13string_searchllPcS_, @function
_Z13string_searchllPcS_:
.LFB3638:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE3638:
	.size	_Z13string_searchllPcS_, .-_Z13string_searchllPcS_
	.globl	_Z16string_search_rrllPcS_icS_
	.type	_Z16string_search_rrllPcS_icS_, @function
_Z16string_search_rrllPcS_icS_:
.LFB3639:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE3639:
	.size	_Z16string_search_rrllPcS_icS_, .-_Z16string_search_rrllPcS_icS_
	.globl	_Z12createVectorll
	.type	_Z12createVectorll, @function
_Z12createVectorll:
.LFB3640:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movq	%rdi, %rbx
	salq	$3, %rdi
	pushq	%rcx
	.cfi_def_cfa_offset 32
	call	malloc@PLT
	xorl	%edx, %edx
.L21:
	cmpq	%rbx, %rdx
	jge	.L24
	movq	%rbp, (%rax,%rdx,8)
	incq	%rdx
	jmp	.L21
.L24:
	popq	%rdx
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3640:
	.size	_Z12createVectorll, .-_Z12createVectorll
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC1:
	.string	"%ld "
.LC2:
	.string	"\n"
	.text
	.globl	_Z10readVectorPll
	.type	_Z10readVectorPll, @function
_Z10readVectorPll:
.LFB3641:
	.cfi_startproc
	endbr64
	pushq	%r13
	.cfi_def_cfa_offset 16
	.cfi_offset 13, -16
	leaq	.LC1(%rip), %r13
	pushq	%r12
	.cfi_def_cfa_offset 24
	.cfi_offset 12, -24
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 32
	.cfi_offset 6, -32
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 40
	.cfi_offset 3, -40
	xorl	%ebx, %ebx
	pushq	%rcx
	.cfi_def_cfa_offset 48
.L26:
	cmpq	%rbp, %rbx
	jge	.L29
	movq	(%r12,%rbx,8), %rdx
	movq	%r13, %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	incq	%rbx
	call	__printf_chk@PLT
	jmp	.L26
.L29:
	popq	%rdx
	.cfi_def_cfa_offset 40
	leaq	.LC2(%rip), %rsi
	popq	%rbx
	.cfi_def_cfa_offset 32
	movl	$2, %edi
	popq	%rbp
	.cfi_def_cfa_offset 24
	xorl	%eax, %eax
	popq	%r12
	.cfi_def_cfa_offset 16
	popq	%r13
	.cfi_def_cfa_offset 8
	jmp	__printf_chk@PLT
	.cfi_endproc
.LFE3641:
	.size	_Z10readVectorPll, .-_Z10readVectorPll
	.globl	_Z9sumVectorPll
	.type	_Z9sumVectorPll, @function
_Z9sumVectorPll:
.LFB3642:
	.cfi_startproc
	endbr64
	xorl	%eax, %eax
	xorl	%edx, %edx
.L31:
	cmpq	%rsi, %rax
	jge	.L33
	addq	(%rdi,%rax,8), %rdx
	incq	%rax
	jmp	.L31
.L33:
	movq	%rdx, %rax
	ret
	.cfi_endproc
.LFE3642:
	.size	_Z9sumVectorPll, .-_Z9sumVectorPll
	.globl	_Z48__device_stub__Z17cuda_stringsearchllPcS_PiiPlS_llPcS_PiiPlS_
	.type	_Z48__device_stub__Z17cuda_stringsearchllPcS_PiiPlS_llPcS_PiiPlS_, @function
_Z48__device_stub__Z17cuda_stringsearchllPcS_PiiPlS_llPcS_PiiPlS_:
.LFB3668:
	.cfi_startproc
	endbr64
	subq	$216, %rsp
	.cfi_def_cfa_offset 224
	movq	224(%rsp), %rax
	movq	%rdi, 56(%rsp)
	leaq	88(%rsp), %rdi
	movq	%rsi, 48(%rsp)
	leaq	100(%rsp), %rsi
	movq	%rax, 8(%rsp)
	movq	232(%rsp), %rax
	movq	%rdx, 40(%rsp)
	leaq	72(%rsp), %rdx
	movq	%rcx, 32(%rsp)
	leaq	80(%rsp), %rcx
	movq	%r8, 24(%rsp)
	movl	%r9d, 20(%rsp)
	movq	%rax, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 200(%rsp)
	xorl	%eax, %eax
	leaq	56(%rsp), %rax
	movl	$1, 96(%rsp)
	movq	%rax, 136(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	20(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 184(%rsp)
	movq	%rsp, %rax
	movq	%rax, 192(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 88(%rsp)
	movq	%rax, 100(%rsp)
	movl	$1, 108(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L34
	pushq	80(%rsp)
	.cfi_def_cfa_offset 232
	leaq	_Z17cuda_stringsearchllPcS_PiiPlS_(%rip), %rdi
	pushq	80(%rsp)
	.cfi_def_cfa_offset 240
	movq	116(%rsp), %rcx
	movl	124(%rsp), %r8d
	movq	104(%rsp), %rsi
	movl	112(%rsp), %edx
	leaq	152(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L34:
	movq	200(%rsp), %rax
	subq	%fs:40, %rax
	je	.L36
	call	__stack_chk_fail@PLT
.L36:
	addq	$216, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3668:
	.size	_Z48__device_stub__Z17cuda_stringsearchllPcS_PiiPlS_llPcS_PiiPlS_, .-_Z48__device_stub__Z17cuda_stringsearchllPcS_PiiPlS_llPcS_PiiPlS_
	.globl	_Z17cuda_stringsearchllPcS_PiiPlS_
	.type	_Z17cuda_stringsearchllPcS_PiiPlS_, @function
_Z17cuda_stringsearchllPcS_PiiPlS_:
.LFB3669:
	.cfi_startproc
	endbr64
	jmp	_Z48__device_stub__Z17cuda_stringsearchllPcS_PiiPlS_llPcS_PiiPlS_
	.cfi_endproc
.LFE3669:
	.size	_Z17cuda_stringsearchllPcS_PiiPlS_, .-_Z17cuda_stringsearchllPcS_PiiPlS_
	.section	.rodata.str1.1
.LC3:
	.string	"r"
.LC4:
	.string	"File error"
.LC5:
	.string	"Buffer index size %lu \n"
.LC6:
	.string	"Question Buffer index size %lu \n"
.LC7:
	.string	"This is question file --------\n"
.LC8:
	.string	"Question max length : %d\n"
.LC9:
	.string	"Question elements count : %d\n"
.LC10:
	.string	"Overflow String size : %d\n"
.LC11:
	.string	"a"
.LC12:
	.string	"This will send buffer start at %ld to %ld of all %ld\n"
.LC13:
	.string	"HOST :: Finish iteration %d at question : %s  temporary founded %ld\n\n"
.LC14:
	.string	"/home/ubuntu/Datasets/stackv2/train-structured/XWinterVarit/LearnParallel/master/cuda3.cu"
.LC15:
	.string	"---------saving change buffered----------------------------------------------------------------\n"
.LC16:
	.string	"%c"
.LC17:
	.string	"-------------------------------------------------------------------------------------------\n"
.LC18:
	.string	"element at : %d is : %s finally founded : %lu\n"
.LC19:
	.string	"\nestimate using time : %.2f\n"
	.section	.text.startup,"ax",@progbits
	.globl	main
	.type	main, @function
main:
.LFB3643:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	xorl	%edi, %edi
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	leaq	.LC3(%rip), %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	subq	$168, %rsp
	.cfi_def_cfa_offset 224
	movq	%fs:40, %rax
	movq	%rax, 152(%rsp)
	xorl	%eax, %eax
	call	time@PLT
	movq	%rbp, %rsi
	leaq	_ZL8FileName(%rip), %rdi
	movq	%rax, 72(%rsp)
	call	fopen@PLT
	movq	%rax, 8(%rsp)
	testq	%rax, %rax
	je	.L73
	movq	8(%rsp), %rdi
	movl	$2, %edx
	xorl	%esi, %esi
	call	fseek@PLT
	movq	8(%rsp), %rdi
	call	ftell@PLT
	movq	8(%rsp), %rdi
	movq	%rax, %rbx
	call	rewind@PLT
	movq	%rbx, %rdx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC5(%rip), %rsi
	call	__printf_chk@PLT
	movq	%rbp, %rsi
	leaq	_ZL16QuestionFileName(%rip), %rdi
	call	fopen@PLT
	movq	%rax, %r12
	testq	%rax, %rax
	je	.L73
	movl	$2, %edx
	xorl	%esi, %esi
	movq	%rax, %rdi
	xorl	%r14d, %r14d
	call	fseek@PLT
	movq	%r12, %rdi
	call	ftell@PLT
	movq	%r12, %rdi
	movq	%rax, %r13
	call	rewind@PLT
	movq	%r13, %rdx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC6(%rip), %rsi
	call	__printf_chk@PLT
	movq	%r13, %rdi
	call	malloc@PLT
	movq	%r12, %r8
	movq	%r13, %rcx
	movl	$1, %edx
	movq	%r13, %rsi
	movq	%rax, %rdi
	movq	%rax, %rbp
	xorl	%r13d, %r13d
	call	__fread_chk@PLT
	leaq	.LC7(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	$16384, %edi
	call	malloc@PLT
	movl	$16384, %edi
	movq	%rax, 48(%rsp)
	call	malloc@PLT
	xorl	%edi, %edi
	movl	%edi, 36(%rsp)
	movq	%rax, %r12
	movl	%edi, 32(%rsp)
.L42:
	movq	%rbp, %rdi
	call	strlen@PLT
	cmpq	%r13, %rax
	jb	.L74
	incq	%r13
	movb	-1(%rbp,%r13), %al
	testb	%al, %al
	je	.L61
	cmpb	$10, %al
	jne	.L42
.L61:
	movl	$30720, %edi
	incl	36(%rsp)
	movslq	36(%rsp), %r15
	call	malloc@PLT
	movq	%r13, %rdx
	leaq	0(%rbp,%r14), %rsi
	movl	$30720, %ecx
	movq	%rax, %rdi
	movq	48(%rsp), %rax
	subq	%r14, %rdx
	movq	%r13, %r14
	decq	%rdx
	movq	%rdi, (%rax,%r15,8)
	call	__memcpy_chk@PLT
	xorl	%esi, %esi
	movq	%rsi, (%r12,%r15,8)
	movq	%rax, %rdi
	call	strlen@PLT
	movslq	32(%rsp), %rdx
	cmpq	%rax, %rdx
	cmovnb	32(%rsp), %eax
	movl	%eax, 32(%rsp)
	jmp	.L42
.L74:
	movl	32(%rsp), %edx
	leaq	.LC8(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	36(%rsp), %edx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC9(%rip), %rsi
	call	__printf_chk@PLT
	movq	%rbp, %rdi
	call	free@PLT
	leaq	.LC7(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movl	32(%rsp), %edx
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC10(%rip), %rsi
	call	__printf_chk@PLT
	leaq	.LC11(%rip), %rsi
	leaq	_ZL14OutputFileName(%rip), %rdi
	call	fopen@PLT
	movq	%rax, 56(%rsp)
	testq	%rax, %rax
	jne	.L46
.L73:
	movq	stderr(%rip), %rsi
	leaq	.LC4(%rip), %rdi
	call	fputs@PLT
	movl	$1, %edi
	call	exit@PLT
.L46:
	xorl	%ecx, %ecx
	movslq	32(%rsp), %r14
	movl	%ecx, 84(%rsp)
	movl	%ecx, 64(%rsp)
.L55:
	leaq	300000000(%r14), %rax
	leaq	299999999(%r14), %r13
	movl	$1, %ebp
	movq	%rax, %rdi
	movq	%rax, 16(%rsp)
	call	malloc@PLT
	movslq	64(%rsp), %rsi
	movq	8(%rsp), %rdi
	movl	$1, %edx
	movq	%rax, 24(%rsp)
	call	fseek@PLT
	movl	32(%rsp), %ecx
	movq	8(%rsp), %r8
	movl	$1, %eax
	movq	16(%rsp), %rsi
	movq	24(%rsp), %rdi
	movl	$1, %edx
	subl	%ecx, %eax
	movq	%r13, %rcx
	movl	%eax, 64(%rsp)
	call	__fread_chk@PLT
	cmpq	$300000000, %rbx
	movq	%rbx, %r8
	leaq	.LC12(%rip), %rsi
	cmovle	%rbx, %r13
	movl	$2, %edi
	xorl	%edx, %edx
	xorl	%eax, %eax
	movq	%r13, %rcx
	call	__printf_chk@PLT
	movq	16(%rsp), %rsi
	leaq	88(%rsp), %rdi
	call	cudaMalloc@PLT
	movq	16(%rsp), %rsi
	leaq	96(%rsp), %rdi
	call	cudaMalloc@PLT
	leaq	104(%rsp), %rdi
	movl	$4, %esi
	call	cudaMalloc@PLT
	movq	16(%rsp), %rdx
	movq	24(%rsp), %rsi
	movl	$1, %ecx
	movq	88(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	16(%rsp), %rdx
	movq	24(%rsp), %rsi
	movl	$1, %ecx
	movq	96(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	104(%rsp), %rdi
	leaq	84(%rsp), %rsi
	movl	$1, %ecx
	movl	$4, %edx
	call	cudaMemcpy@PLT
.L48:
	movl	%ebp, 68(%rsp)
	cmpl	%ebp, 36(%rsp)
	jle	.L75
	xorl	%esi, %esi
	movl	$30720, %edi
	call	_Z12createVectorll
	leaq	120(%rsp), %rdi
	movl	$245760, %esi
	movq	%rax, 40(%rsp)
	call	cudaMalloc@PLT
	leaq	112(%rsp), %rdi
	movl	$2048, %esi
	call	cudaMalloc@PLT
	movq	40(%rsp), %rsi
	movq	120(%rsp), %rdi
	movl	$1, %ecx
	movl	$245760, %edx
	call	cudaMemcpy@PLT
	movq	48(%rsp), %rax
	movq	112(%rsp), %rdi
	movl	$1, %ecx
	movl	$2048, %edx
	movq	(%rax,%rbp,8), %r15
	movq	%r15, %rsi
	call	cudaMemcpy@PLT
	movl	$4194305, %edx
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$2147483663, %edi
	salq	$10, %rdx
	movl	$1, %ecx
	movl	$1, %esi
	addq	%rdi, %rdi
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L49
	pushq	96(%rsp)
	.cfi_def_cfa_offset 232
	movq	%r13, %rsi
	xorl	%edi, %edi
	pushq	128(%rsp)
	.cfi_def_cfa_offset 240
	movl	48(%rsp), %r9d
	movq	120(%rsp), %r8
	movq	104(%rsp), %rcx
	movq	128(%rsp), %rdx
	call	_Z48__device_stub__Z17cuda_stringsearchllPcS_PiiPlS_llPcS_PiiPlS_
	popq	%rax
	.cfi_def_cfa_offset 232
	popq	%rdx
	.cfi_def_cfa_offset 224
.L49:
	call	cudaDeviceSynchronize@PLT
	movq	120(%rsp), %rsi
	movq	40(%rsp), %rdi
	movl	$245760, %edx
	movl	$2, %ecx
	call	cudaMemcpy@PLT
	movq	120(%rsp), %rdi
	call	cudaFree@PLT
	movq	112(%rsp), %rdi
	call	cudaFree@PLT
	movq	40(%rsp), %rdi
	movl	$30720, %esi
	movq	%r15, %rcx
	call	_Z9sumVectorPll
	addq	%rax, (%r12,%rbp,8)
	movl	68(%rsp), %edx
	incq	%rbp
	movq	%rax, %r8
	leaq	.LC13(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	call	__printf_chk@PLT
	movq	40(%rsp), %rdi
	call	free@PLT
	jmp	.L48
.L75:
	call	cudaPeekAtLastError@PLT
	leaq	.LC14(%rip), %rbp
	movl	$1, %ecx
	movl	$354, %edx
	movq	%rbp, %rsi
	movl	%eax, %edi
	leaq	.LC16(%rip), %r13
	call	_Z9gpuAssert9cudaErrorPKcib
	call	cudaDeviceSynchronize@PLT
	movq	%rbp, %rsi
	movl	$1, %ecx
	movl	$355, %edx
	movl	%eax, %edi
	movl	$299999999, %ebp
	call	_Z9gpuAssert9cudaErrorPKcib
	movq	16(%rsp), %rdx
	movq	96(%rsp), %rsi
	movl	$2, %ecx
	movq	24(%rsp), %rdi
	call	cudaMemcpy@PLT
	movq	88(%rsp), %rdi
	call	cudaFree@PLT
	movq	104(%rsp), %rdi
	call	cudaFree@PLT
	movq	96(%rsp), %rdi
	call	cudaFree@PLT
	xorl	%eax, %eax
	movl	$2, %edi
	leaq	.LC15(%rip), %rsi
	call	__printf_chk@PLT
	cmpq	%rbp, %rbx
	cmovle	%rbx, %rbp
	xorl	%r15d, %r15d
.L51:
	cmpq	%r15, %rbp
	jl	.L76
	movq	24(%rsp), %rax
	movsbl	(%rax,%r15), %ecx
	cmpb	$36, %cl
	je	.L52
	testb	%cl, %cl
	je	.L52
	movq	56(%rsp), %rdi
	movq	%r13, %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L52:
	incq	%r15
	jmp	.L51
.L76:
	leaq	.LC2(%rip), %rsi
	movl	$2, %edi
	xorl	%eax, %eax
	subq	$300000000, %rbx
	call	__printf_chk@PLT
	movl	$2, %edi
	xorl	%eax, %eax
	leaq	.LC17(%rip), %rsi
	call	__printf_chk@PLT
	movq	24(%rsp), %rdi
	call	free@PLT
	testq	%rbx, %rbx
	jg	.L55
	movl	$1, %ebp
	leaq	.LC18(%rip), %rbx
.L56:
	movl	%ebp, %edx
	cmpl	%ebp, 36(%rsp)
	jl	.L77
	movq	48(%rsp), %rax
	movq	(%r12,%rbp,8), %r8
	movq	%rbx, %rsi
	movl	$2, %edi
	movq	(%rax,%rbp,8), %rcx
	xorl	%eax, %eax
	incq	%rbp
	call	__printf_chk@PLT
	jmp	.L56
.L77:
	movq	8(%rsp), %rdi
	call	fclose@PLT
	movq	56(%rsp), %rdi
	call	fclose@PLT
	xorl	%edi, %edi
	call	time@PLT
	movq	72(%rsp), %rbx
	movl	$2, %edi
	leaq	.LC19(%rip), %rsi
	subq	%rbx, %rax
	cvtsi2sdq	%rax, %xmm0
	movb	$1, %al
	call	__printf_chk@PLT
	movq	152(%rsp), %rax
	subq	%fs:40, %rax
	je	.L58
	call	__stack_chk_fail@PLT
.L58:
	addq	$168, %rsp
	.cfi_def_cfa_offset 56
	xorl	%eax, %eax
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3643:
	.size	main, .-main
	.section	.rodata.str1.1
.LC20:
	.string	"_Z17cuda_stringsearchllPcS_PiiPlS_"
	.section	.text.startup
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB3671:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC20(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z17cuda_stringsearchllPcS_PiiPlS_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE3671:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.globl	WordCount
	.bss
	.align 4
	.type	WordCount, @object
	.size	WordCount, 4
WordCount:
	.zero	4
	.section	.rodata
	.type	_ZL16QuestionFileName, @object
	.size	_ZL16QuestionFileName, 6
_ZL16QuestionFileName:
	.string	"line1"
	.type	_ZL14OutputFileName, @object
	.size	_ZL14OutputFileName, 6
_ZL14OutputFileName:
	.string	"v.txt"
	.type	_ZL8FileName, @object
	.size	_ZL8FileName, 6
_ZL8FileName:
	.string	"01.nt"
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:


#include <hip/hip_runtime.h>
#include <iostream>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>

#define BLOCKDIM 1024
/*
#include <thrust/device_vector.h>
#include <thrust/reduce.h>
#include <thrust/functional.h>
#include <thrust/inner_product.h>
using namespace thrust;
using namespace thrust::placeholders;
*/
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(hipError_t code, const char *file, int line, bool abort=true)
{
    if (code != hipSuccess)
    {
        fprintf(stderr,"GPUassert: %s %s %d\n", hipGetErrorString(code), file, line);
        if (abort) exit(code);
    }
}

/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
///////////////////////////////Important Configuration/////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////

const char FileName[] = "01.nt"; // the file must have
const char OutputFileName[] = "v.txt"; // the file must have, but must be empty
const char QuestionFileName[] = "line1"; // the file must have


// Threadchucksize mean "Size" in byte of part of the file that will be sent to the gpu memory
// not worry about setting size higher than real file, because the program will auto adjust it
// Example , if your gpu memory limit at 500 MB but your file is large such as 5 GB , you should set threadchucksize to 500 * 1,000,000byte = 500,000,000
// and your part of file will be sent 5000 MB/500 MB = 10 times
// if you sent and found some memory error, it can be the os or other software use vram too, so decrease threadchucksize until no error, such as from 500,000,000 change to 200,000,000
// suggestion, you should set size as beautiful ten-end number as 100000, 20000000, 50000000
const long threadchucksize = 300000000;


// this Blocksize not mean thread per block but mean "size" of data chuck that each thread will compute from the big global data chuck (that locate in gpu)
// warning that allThreadInUse multiply with blocksize must higher than threadchucksize, if not it will incorrect result
// Example : if threadchucksize = 100,000,000 (aka 100MB. chuck of file sent to gpu) and allThreadInUse = 4096 and blocksize = 25000
// then you must check that 25000 * 4096 > 100000000 ? which is 102,400,000 > 100,000,000 so it true and can be use
// suggestion, you should set size as threadchucksize % blocksize = 0, it will be bug free.
const long blocksize = 10000;


const int NumberOfComputeBlock = 30; // aka gridsize
const int NumberOfThreadsPerBlock = 1024; //(rely on your gpu spec)

// allThreadInUse is all concurrent thread that run in the gpu,
// it can be higher than physical cuda core on gpu, because gpu can queue it and make you feel like it concurrent
// but if allThreadInUse is much higher, the answer vector that collect answer from each thread will be larger. so threadoff
const int allThreadInUse = NumberOfThreadsPerBlock * NumberOfComputeBlock;
const long sizeofAnswerVector = allThreadInUse;


const int sizeofQuestionArray = 2048; // maximum list size of question, such as = 2048 mean this program support maximum 2048 question
const int sizeofMaximumQuestionWord = 2048; // maximum string size of each question, such as = 2048 mean each question can't be larger than 2048 byte


/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////
/////////////////////////////////////////////////////////////////////////////////////



int WordCount;
__device__ size_t d_strlen (const char *str)
{
    return (*str) ? d_strlen(++str) + 1 : 0;
}
__device__ int d_strncmp(const char *ptr0, const char *ptr1, size_t len)
{
    while(len--)
        if(*ptr0++!=*ptr1++)
            return *(unsigned char*)(ptr0 - 1) - *(unsigned char*)(ptr1 - 1);
    return 0;
}
__device__ unsigned int string_search(long start, long end, char* target, char *buffer) {
    unsigned int i;
    unsigned int found=0;
    //printf("Receiveing : target : %s  buffer : %s\n", target, buffer);
    //printf("My Buffer: ");
    //for (i=start;i <= end; i++)
    //    printf("%c", *(buffer+i));
    //printf("\n");
    for (i=start;i <= end; i++) {
        int t = d_strncmp(&buffer[i], target, d_strlen(target));
        //printf("t dkmfdsfdfdspfdsfpodsfjkdpsof: %d \n", t);
        if (t == 0) {
            //if (i <= overflowRegion)
            found++;
        }
    }
    //printf("overflowRegion : %d\n", overflowRegion);
    //printf("Receiveing Found : %d\n", found);
    return found;
}
__device__ unsigned int string_search_rr(long start, long end, char* target, char *buffer,int overflowStringSize, char options, char *changebuffer) {

    unsigned int i;
    unsigned int found=0;

    for (i=start;i <= end  ; i++) {
        int t = d_strncmp(&buffer[i], target, d_strlen(target));
        if (t == 0 ) {
            //if (i <= overflowRegion)
            found++;

            for (int j = i; j < i + d_strlen(target); ++j) {
                //printf("change at j : %d i : %d\n", j,i);
                *(changebuffer+j) = '$';
            }

        } else {
	    // if you want to search word in any file (not only file that contain one line per word. please comment this block)
	    //
	    int j = 0;
 	    for (j = i; j <= end; j++) {
	    	if (*(buffer+j) == '\n')
	 	    break;
	    }
	    i = j;
	    //
	    //-----------------------------------------------------------------------------------------------------------------
	}

    }

    return found;
}
__global__ void cuda_stringsearch (long bufferstart, long bufferend, char* target, char* buffer, int* allcount, int overflowStringSize, long *answerVector, char* changebuffer) {
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    //printf("hello from thread %d\n", index);
/*  
  if (index == 1) {
        printf("GPU KERNEL :: Hello from threads %d  Given word %s  size %d \n", index, target, d_strlen(target));
    }
*/
    //printf("Hello from threads %d  Given word %s\n", index, target);

   // long blocksize = 500/*50000*/;
    long extendblocksize = blocksize + overflowStringSize - 2;
    long startpoint = index * blocksize;
    long endpoint = startpoint + blocksize  - 1;

    if (startpoint <= bufferend) {
        if (endpoint > bufferend)
            endpoint = bufferend;

        //int count = 10;
        int count = string_search_rr(startpoint, endpoint,target, buffer, overflowStringSize, 'd', changebuffer);
        //printf("threads %d count %d  getting data :  startpoint %ld  endpoint %ld  overflowStringSize %d\n", index,count,startpoint, endpoint, overflowStringSize);
        *(answerVector + index) = count;
        //*allcount += count;
    }

}

long* createVector (long size, long inivalue) {
    long* vector = (long*) malloc(sizeof(long)*size);
    for (long i = 0; i < size; ++i) {
        vector[i] = inivalue;
    }
    return vector;
}
void readVector (long* vector, long size) {
    for (long i = 0; i < size; ++i) {
        printf("%ld ",vector[i]);
    }
    printf("\n");
}
long sumVector (long* vector, long size) {
    long sum = 0;
    for (long i = 0; i < size; ++i) {
        sum+= vector[i];
    }
    return sum;
}
int main(int argc, char **argv) {

    time_t timestart = time(NULL);

    FILE * pFile;
    long lSize;
    pFile = fopen ( FileName , "r" );
    if (pFile==NULL) {fputs ("File error",stderr); exit (1);}
    fseek (pFile , 0 , SEEK_END);
    lSize = ftell (pFile);
    rewind (pFile);
    long BufferSize = sizeof(char)*lSize;
    printf("Buffer index size %lu \n",BufferSize);

    int reverseoffset = 0;
    /* create threads */
    long endpoint = 0,startpoint = 0;




    FILE * questionFile;
    long lSizeQ;
    questionFile = fopen(QuestionFileName, "r");
    if (questionFile == NULL) {fputs ("File error", stderr); exit(1);}
    fseek(questionFile, 0, SEEK_END);
    lSizeQ = ftell(questionFile);
    rewind(questionFile);
    long QuestionBufferSize = sizeof(char)*lSizeQ;
    printf("Question Buffer index size %lu \n", QuestionBufferSize);
    char *Question_Buffer = (char*) malloc (lSizeQ);
    fread(Question_Buffer, 1, QuestionBufferSize, questionFile);
    printf("This is question file --------\n");

    //printf("%s\n", Question_Buffer);

    long start = 0, end = 0;
    int Question_maxLength = 0;
    char** questionArray = (char**) malloc(sizeof(char*)*sizeofQuestionArray);
    long* questionAnswer = (long*) malloc(sizeof(long)*sizeofQuestionArray); // found word list

    int questionCount = 0;
    for (int j = 0; j <= strlen(Question_Buffer); ++j) {
        end++;
        if (*(Question_Buffer+j) == '\n' || *(Question_Buffer+j) == '\0') {
            questionCount++;
            //piece = (char*) malloc(sizeof(char)*2048);
            //memcpy(piece, (Question_Buffer+start), end - start - 1);
            *(questionArray+questionCount) = (char*) malloc(sizeof(char)*sizeofAnswerVector);
            memcpy(*(questionArray+questionCount), (Question_Buffer+start), end - start - 1);
            *(questionAnswer+questionCount) = 0; /* each question start founded = zero */               /*(long) strlen(*(questionArray+questionCount)) this commented code use to check if for loop work!*/
            if (strlen(*(questionArray+questionCount)) > Question_maxLength)
                Question_maxLength = strlen(*(questionArray+questionCount));
            //printf("print piece %s|||\n", *(questionArray+questionCount));
            //printf("piece length : %lu \n", strlen(*(questionArray+questionCount)));
            start = end;

        }
    }
/*
    for (int k = 1; k <= questionCount; ++k) {
        printf("element at : %d is : %s value is : %lu  length is %zu\n",  k, *(questionArray+k), *(questionAnswer+k), strlen(*(questionArray+k)));
    }
*/
    printf("Question max length : %d\n", Question_maxLength);
    printf("Question elements count : %d\n", questionCount);
    free(Question_Buffer);
    printf("This is question file --------\n");

    int overflowStringSize = Question_maxLength /*- 1*/;
    printf("Overflow String size : %d\n", overflowStringSize);



    FILE * outputFile;
    long lSize2;
    outputFile = fopen(OutputFileName, "a");
    if (outputFile==NULL) {fputs ("File error",stderr); exit (1);}


    int count = 0;
    int* countPTR = &count;
    int overflowRegion = threadchucksize - 1;
    while (1){
        //printf("precount  all count %d\n", count);

        char *buffer;
        startpoint = 0;
        endpoint = threadchucksize + overflowStringSize - 1;
        buffer = (char*) malloc (sizeof(char)*(threadchucksize + overflowStringSize));
        fseek (pFile , reverseoffset , SEEK_CUR);
        reverseoffset = -1 * (overflowStringSize - 1);

        fread (buffer,1,endpoint,pFile);
        if (BufferSize <= threadchucksize)
            endpoint = BufferSize;
        printf("This will send buffer start at %ld to %ld of all %ld\n", startpoint, endpoint, BufferSize);

        //int j = 0;
        //count += string_search(startpoint, endpoint, DefineWord, buffer);
        //printf("%s||| count : %d\n", buffer, count);
        //printf("-------\n");




        char *dev_buffer;
        char *dev_changebuffer;
        int *dev_countPTR;
        //char *dev_defineword;
        //long *dev_answerVector;

        hipMalloc((void**)&dev_buffer, sizeof(char)*(threadchucksize + overflowStringSize));
        hipMalloc((void**)&dev_changebuffer, sizeof(char)*(threadchucksize + overflowStringSize));
        hipMalloc((void**)&dev_countPTR, sizeof(int));

        hipMemcpy(dev_buffer, buffer, sizeof(char)*(threadchucksize + overflowStringSize), hipMemcpyHostToDevice);
        hipMemcpy(dev_changebuffer, buffer, sizeof(char)*(threadchucksize + overflowStringSize), hipMemcpyHostToDevice);

        hipMemcpy(dev_countPTR, countPTR, sizeof(int),hipMemcpyHostToDevice);

        for (int question = 1; question <= questionCount - 1; ++question) {
        //int question = 1;
            long size_answerVector = sizeofAnswerVector;
            long* answerVector = createVector(size_answerVector,0);
	    char *dev_defineword;
	    long *dev_answerVector;

//	    printf("HOST :: starting iteration %d at question : %s  string length : %zu\n",question, *(questionArray+question),  strlen(*(questionArray+question)));
/*		for (int d = 0; d < strlen(*(questionArray+question)); d++) {
			printf("%c",*(*(questionArray+question)+d));
		}
		printf("\nend test \n");*/
            hipMalloc((void**)&dev_answerVector, sizeof(long)*size_answerVector);
            hipMalloc((void**)&dev_defineword, /*sizeof(char)**/ sizeofMaximumQuestionWord/*strlen(*(questionArray+question))*/);

            hipMemcpy(dev_answerVector, answerVector, sizeof(long)*size_answerVector, hipMemcpyHostToDevice);
            hipMemcpy(dev_defineword, *(questionArray+question), /*sizeof(char)**/ sizeofMaximumQuestionWord /*strlen(*(questionArray+question))*/, hipMemcpyHostToDevice);
            //printf("iteration at question : %s\n", *(questionArray+question));

            cuda_stringsearch<<<NumberOfComputeBlock,NumberOfThreadsPerBlock>>>(startpoint, endpoint, dev_defineword, dev_buffer, dev_countPTR, overflowStringSize, dev_answerVector, dev_changebuffer);
            hipDeviceSynchronize();
            hipMemcpy (answerVector, dev_answerVector, sizeof(long)*size_answerVector, hipMemcpyDeviceToHost);

            hipFree(dev_answerVector);
            hipFree(dev_defineword);
	    
            //readVector(answerVector, size_answerVector); //uncomment this to diagnostic answer vector matrix
            long iterationsum =  sumVector(answerVector, size_answerVector);

            *(questionAnswer+question) += iterationsum;

            printf("HOST :: Finish iteration %d at question : %s  temporary founded %ld\n\n",question, *(questionArray+question),  iterationsum);
            free(answerVector);
        }


        gpuErrchk( hipPeekAtLastError() );
        gpuErrchk( hipDeviceSynchronize() );

        hipMemcpy (buffer, dev_changebuffer,sizeof(char)*(threadchucksize + overflowStringSize),hipMemcpyDeviceToHost);
        //cudaMemcpy (countPTR, dev_countPTR, sizeof(int), cudaMemcpyDeviceToHost);
        hipFree(dev_buffer); hipFree(dev_countPTR); hipFree(dev_changebuffer);

        printf("---------saving change buffered----------------------------------------------------------------\n");
        endpoint = threadchucksize - 1;
        if (endpoint > BufferSize)
            endpoint = BufferSize;
        //printf("startpoint : %ld     endpoint : %ld   BufferSize : %ld \n", startpoint, endpoint, BufferSize);
        for (int i=startpoint;i <= endpoint; i++) {
            if (*(buffer + i) != '\0' && *(buffer + i) != '$' )
                fprintf(outputFile, "%c", *(buffer + i));
            //printf("%c", *(buffer + i));
        }
        printf("\n");

        printf("-------------------------------------------------------------------------------------------\n");

        //fprintf(outputFile, "%s",buffer);

        BufferSize = BufferSize - threadchucksize;

        free(buffer);
        if (BufferSize <= 0)
            break;
    }
    for (int k = 1; k <= questionCount; ++k) {
        printf("element at : %d is : %s finally founded : %lu\n",  k, *(questionArray+k), *(questionAnswer+k));
    }
    fclose (pFile);
    fclose (outputFile);

    printf("\nestimate using time : %.2f\n", (double)(time(NULL) - timestart));

    return EXIT_SUCCESS;
}

	.text
	.file	"cuda3.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.globl	_Z32__device_stub__cuda_stringsearchllPcS_PiiPlS_ # -- Begin function _Z32__device_stub__cuda_stringsearchllPcS_PiiPlS_
	.type	_Z32__device_stub__cuda_stringsearchllPcS_PiiPlS_,@function
_Z32__device_stub__cuda_stringsearchllPcS_PiiPlS_: # @_Z32__device_stub__cuda_stringsearchllPcS_PiiPlS_
	.cfi_startproc
# %bb.0:
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%r13
	.cfi_def_cfa_offset 32
	pushq	%r12
	.cfi_def_cfa_offset 40
	pushq	%rbx
	.cfi_def_cfa_offset 48
	subq	$160, %rsp
	.cfi_def_cfa_offset 208
	.cfi_offset %rbx, -48
	.cfi_offset %r12, -40
	.cfi_offset %r13, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	leaq	56(%rsp), %rax
	movq	%rdi, (%rax)
	leaq	48(%rsp), %rdi
	movq	%rsi, (%rdi)
	leaq	40(%rsp), %rsi
	movq	%rdx, (%rsi)
	leaq	32(%rsp), %rdx
	movq	%rcx, (%rdx)
	leaq	24(%rsp), %rcx
	movq	%r8, (%rcx)
	leaq	4(%rsp), %r8
	movl	%r9d, (%r8)
	leaq	96(%rsp), %rbx
	movq	%rax, (%rbx)
	movq	%rdi, 8(%rbx)
	movq	%rsi, 16(%rbx)
	movq	%rdx, 24(%rbx)
	movq	%rcx, 32(%rbx)
	movq	%r8, 40(%rbx)
	leaq	208(%rsp), %rax
	movq	%rax, 48(%rbx)
	leaq	216(%rsp), %rax
	movq	%rax, 56(%rbx)
	leaq	80(%rsp), %r14
	leaq	64(%rsp), %r15
	leaq	16(%rsp), %r12
	leaq	8(%rsp), %r13
	movq	%r14, %rdi
	movq	%r15, %rsi
	movq	%r12, %rdx
	movq	%r13, %rcx
	callq	__hipPopCallConfiguration
	movq	(%r14), %rsi
	movl	8(%r14), %edx
	movq	(%r15), %rcx
	movl	8(%r15), %r8d
	movl	$_Z17cuda_stringsearchllPcS_PiiPlS_, %edi
	movq	%rbx, %r9
	pushq	(%r13)
	.cfi_adjust_cfa_offset 8
	pushq	(%r12)
	.cfi_adjust_cfa_offset 8
	callq	hipLaunchKernel
	addq	$176, %rsp
	.cfi_adjust_cfa_offset -176
	popq	%rbx
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end0:
	.size	_Z32__device_stub__cuda_stringsearchllPcS_PiiPlS_, .Lfunc_end0-_Z32__device_stub__cuda_stringsearchllPcS_PiiPlS_
	.cfi_endproc
                                        # -- End function
	.globl	_Z12createVectorll              # -- Begin function _Z12createVectorll
	.type	_Z12createVectorll,@function
_Z12createVectorll:                     # @_Z12createVectorll
	.cfi_startproc
# %bb.0:
	pushq	%r14
	.cfi_def_cfa_offset 16
	pushq	%rbx
	.cfi_def_cfa_offset 24
	pushq	%rax
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -24
	.cfi_offset %r14, -16
	movq	%rsi, %rbx
	movq	%rdi, %r14
	leaq	(,%rdi,8), %rdi
	callq	malloc
	testq	%r14, %r14
	jle	.LBB1_3
# %bb.1:                                # %.lr.ph.preheader
	xorl	%ecx, %ecx
.LBB1_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movq	%rbx, (%rax,%rcx,8)
	incq	%rcx
	cmpq	%rcx, %r14
	jne	.LBB1_2
.LBB1_3:                                # %._crit_edge
	addq	$8, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%r14
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z12createVectorll, .Lfunc_end1-_Z12createVectorll
	.cfi_endproc
                                        # -- End function
	.globl	_Z10readVectorPll               # -- Begin function _Z10readVectorPll
	.type	_Z10readVectorPll,@function
_Z10readVectorPll:                      # @_Z10readVectorPll
	.cfi_startproc
# %bb.0:
	testq	%rsi, %rsi
	jle	.LBB2_4
# %bb.1:                                # %.lr.ph.preheader
	pushq	%r15
	.cfi_def_cfa_offset 16
	pushq	%r14
	.cfi_def_cfa_offset 24
	pushq	%rbx
	.cfi_def_cfa_offset 32
	.cfi_offset %rbx, -32
	.cfi_offset %r14, -24
	.cfi_offset %r15, -16
	movq	%rsi, %rbx
	movq	%rdi, %r14
	xorl	%r15d, %r15d
.LBB2_2:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	movq	(%r14,%r15,8), %rsi
	movl	$.L.str, %edi
	xorl	%eax, %eax
	callq	printf
	incq	%r15
	cmpq	%r15, %rbx
	jne	.LBB2_2
# %bb.3:
	popq	%rbx
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	.cfi_restore %rbx
	.cfi_restore %r14
	.cfi_restore %r15
.LBB2_4:                                # %._crit_edge
	movl	$10, %edi
	jmp	putchar@PLT                     # TAILCALL
.Lfunc_end2:
	.size	_Z10readVectorPll, .Lfunc_end2-_Z10readVectorPll
	.cfi_endproc
                                        # -- End function
	.globl	_Z9sumVectorPll                 # -- Begin function _Z9sumVectorPll
	.type	_Z9sumVectorPll,@function
_Z9sumVectorPll:                        # @_Z9sumVectorPll
	.cfi_startproc
# %bb.0:
	testq	%rsi, %rsi
	jle	.LBB3_1
# %bb.3:                                # %.lr.ph.preheader
	xorl	%ecx, %ecx
	xorl	%eax, %eax
.LBB3_4:                                # %.lr.ph
                                        # =>This Inner Loop Header: Depth=1
	addq	(%rdi,%rcx,8), %rax
	incq	%rcx
	cmpq	%rcx, %rsi
	jne	.LBB3_4
# %bb.2:                                # %._crit_edge
	retq
.LBB3_1:
	xorl	%eax, %eax
	retq
.Lfunc_end3:
	.size	_Z9sumVectorPll, .Lfunc_end3-_Z9sumVectorPll
	.cfi_endproc
                                        # -- End function
	.globl	main                            # -- Begin function main
	.type	main,@function
main:                                   # @main
	.cfi_startproc
# %bb.0:
	pushq	%rbp
	.cfi_def_cfa_offset 16
	pushq	%r15
	.cfi_def_cfa_offset 24
	pushq	%r14
	.cfi_def_cfa_offset 32
	pushq	%r13
	.cfi_def_cfa_offset 40
	pushq	%r12
	.cfi_def_cfa_offset 48
	pushq	%rbx
	.cfi_def_cfa_offset 56
	subq	$152, %rsp
	.cfi_def_cfa_offset 208
	.cfi_offset %rbx, -56
	.cfi_offset %r12, -48
	.cfi_offset %r13, -40
	.cfi_offset %r14, -32
	.cfi_offset %r15, -24
	.cfi_offset %rbp, -16
	xorl	%edi, %edi
	callq	time
	movq	%rax, %rbx
	movl	$_ZL8FileName, %edi
	movl	$.L.str.2, %esi
	callq	fopen
	testq	%rax, %rax
	je	.LBB4_32
# %bb.1:
	movq	%rax, %r14
	movq	%rbx, 112(%rsp)                 # 8-byte Spill
	xorl	%ebx, %ebx
	movq	%rax, %rdi
	xorl	%esi, %esi
	movl	$2, %edx
	callq	fseek
	movq	%r14, %rdi
	callq	ftell
	movq	%rax, %r15
	movq	%r14, 96(%rsp)                  # 8-byte Spill
	movq	%r14, %rdi
	callq	rewind
	movl	$.L.str.4, %edi
	movq	%r15, 16(%rsp)                  # 8-byte Spill
	movq	%r15, %rsi
	xorl	%eax, %eax
	callq	printf
	movl	$_ZL16QuestionFileName, %edi
	movl	$.L.str.2, %esi
	callq	fopen
	testq	%rax, %rax
	je	.LBB4_32
# %bb.2:
	movq	%rax, %r14
	movq	%rax, %rdi
	xorl	%esi, %esi
	movl	$2, %edx
	callq	fseek
	movq	%r14, %rdi
	callq	ftell
	movq	%rax, %r15
	movq	%r14, %rdi
	callq	rewind
	movl	$.L.str.5, %edi
	movq	%r15, %rsi
	xorl	%eax, %eax
	callq	printf
	movq	%r15, %rdi
	callq	malloc
	movl	$1, %esi
	movq	%rax, %rdi
	movq	%r15, %rdx
	movq	%rax, %r15
	movq	%r14, %rcx
	callq	fread
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movl	$16384, %edi                    # imm = 0x4000
	callq	malloc
	movq	%rax, 40(%rsp)                  # 8-byte Spill
	movl	$16384, %edi                    # imm = 0x4000
	callq	malloc
	movq	%rax, 32(%rsp)                  # 8-byte Spill
	xorl	%r13d, %r13d
	xorl	%r14d, %r14d
	xorl	%ebp, %ebp
.LBB4_3:                                # =>This Inner Loop Header: Depth=1
	leaq	1(%r14), %r12
	movzbl	(%r15,%r14), %eax
	cmpl	$10, %eax
	je	.LBB4_5
# %bb.4:                                #   in Loop: Header=BB4_3 Depth=1
	testl	%eax, %eax
	jne	.LBB4_6
.LBB4_5:                                #   in Loop: Header=BB4_3 Depth=1
	leal	1(%rbp), %eax
	movl	%eax, 24(%rsp)                  # 4-byte Spill
	movl	$30720, %edi                    # imm = 0x7800
	callq	malloc
	movq	%r13, %rcx
	movq	%r15, %r13
	movq	%rax, %r15
	movq	%rbp, %rax
	movq	%rbx, %rbp
	movslq	%eax, %rbx
	movq	40(%rsp), %rax                  # 8-byte Reload
	movq	%r15, 8(%rax,%rbx,8)
	leaq	(%rcx,%r13), %rsi
	subq	%rcx, %r14
	movq	%r15, %rdi
	movq	%r14, %rdx
	callq	memcpy@PLT
	movq	32(%rsp), %rax                  # 8-byte Reload
	movq	$0, 8(%rax,%rbx,8)
	movq	%rbp, %rbx
	movq	%r15, %rdi
	movq	%r13, %r15
	callq	strlen
	movslq	%ebx, %rbx
	cmpq	%rbx, %rax
	cmoval	%eax, %ebx
	movl	24(%rsp), %eax                  # 4-byte Reload
	movl	%eax, %ebp
	movq	%r12, %r13
.LBB4_6:                                #   in Loop: Header=BB4_3 Depth=1
	movq	%r15, %rdi
	callq	strlen
	movq	%r12, %rcx
	decq	%rcx
	movq	%r12, %r14
	cmpq	%rcx, %rax
	ja	.LBB4_3
# %bb.7:
	movl	$.L.str.7, %edi
	movl	%ebx, %esi
	xorl	%eax, %eax
	callq	printf
	movl	$.L.str.8, %edi
	movl	%ebp, %esi
	xorl	%eax, %eax
	callq	printf
	movq	%r15, %rdi
	callq	free
	movl	$.Lstr.1, %edi
	callq	puts@PLT
	movl	$.L.str.9, %edi
	movl	%ebx, %esi
	xorl	%eax, %eax
	callq	printf
	movl	$_ZL14OutputFileName, %edi
	movl	$.L.str.10, %esi
	callq	fopen
	movq	%rax, 80(%rsp)                  # 8-byte Spill
	testq	%rax, %rax
	je	.LBB4_32
# %bb.8:
	movabsq	$4294967326, %rax               # imm = 0x10000001E
	leaq	76(%rsp), %rcx
	movl	$0, (%rcx)
	movslq	%ebx, %rcx
	leaq	300000000(%rcx), %rdx
	movq	%rdx, 104(%rsp)                 # 8-byte Spill
	addq	$299999999, %rcx                # imm = 0x11E1A2FF
	movq	%rcx, 128(%rsp)                 # 8-byte Spill
	movl	$1, %ecx
	movq	%rbx, 136(%rsp)                 # 8-byte Spill
	subl	%ebx, %ecx
	movl	%ecx, 72(%rsp)                  # 4-byte Spill
	movq	%rbp, 88(%rsp)                  # 8-byte Spill
	movl	%ebp, %ecx
	movq	%rcx, 120(%rsp)                 # 8-byte Spill
	addq	$994, %rax                      # imm = 0x3E2
	movq	%rax, 24(%rsp)                  # 8-byte Spill
	xorl	%ebp, %ebp
	movq	96(%rsp), %r13                  # 8-byte Reload
	movq	16(%rsp), %r15                  # 8-byte Reload
.LBB4_9:                                # =>This Loop Header: Depth=1
                                        #     Child Loop BB4_11 Depth 2
                                        #       Child Loop BB4_14 Depth 3
                                        #     Child Loop BB4_23 Depth 2
	movl	$299999999, %eax                # imm = 0x11E1A2FF
	cmpq	%rax, %r15
	movl	$299999999, %r12d               # imm = 0x11E1A2FF
	cmovlq	%r15, %r12
	testq	%r12, %r12
	movl	$0, %eax
	cmovleq	%rax, %r12
	movq	104(%rsp), %r14                 # 8-byte Reload
	movq	%r14, %rdi
	callq	malloc
	movq	%rax, %rbx
	movslq	%ebp, %rsi
	movq	%r13, %rdi
	movl	$1, %edx
	callq	fseek
	movl	$1, %esi
	movq	%rbx, %rdi
	movq	%r13, %rcx
	movq	128(%rsp), %r13                 # 8-byte Reload
	movq	%r13, %rdx
	callq	fread
	cmpq	$300000001, %r15                # imm = 0x11E1A301
	cmovlq	%r15, %r13
	movl	$.L.str.11, %edi
	xorl	%esi, %esi
	movq	%r13, 144(%rsp)                 # 8-byte Spill
	movq	%r13, %rdx
	movq	%r15, 16(%rsp)                  # 8-byte Spill
	movq	%r15, %rcx
	xorl	%eax, %eax
	callq	printf
	leaq	64(%rsp), %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	leaq	8(%rsp), %rdi
	movq	%r14, %rsi
	callq	hipMalloc
	movl	$4, %esi
	leaq	56(%rsp), %rdi
	callq	hipMalloc
	movq	64(%rsp), %rdi
	movq	%rbx, %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	8(%rsp), %rdi
	movq	%rbx, %rsi
	movq	%r14, %rdx
	movl	$1, %ecx
	callq	hipMemcpy
	movq	56(%rsp), %rdi
	movl	$4, %edx
	leaq	76(%rsp), %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	cmpl	$2, 88(%rsp)                    # 4-byte Folded Reload
	movq	120(%rsp), %r14                 # 8-byte Reload
	jl	.LBB4_16
# %bb.10:                               # %.lr.ph.preheader
                                        #   in Loop: Header=BB4_9 Depth=1
	movl	$1, %r15d
.LBB4_11:                               # %.lr.ph
                                        #   Parent Loop BB4_9 Depth=1
                                        # =>  This Loop Header: Depth=2
                                        #       Child Loop BB4_14 Depth 3
	movl	$1, %edi
	movl	$245760, %esi                   # imm = 0x3C000
	callq	calloc@PLT
	movq	%rax, %r13
	movl	$245760, %esi                   # imm = 0x3C000
	movq	%rsp, %rdi
	callq	hipMalloc
	movl	$2048, %esi                     # imm = 0x800
	leaq	48(%rsp), %rdi
	callq	hipMalloc
	movq	(%rsp), %rdi
	movl	$245760, %edx                   # imm = 0x3C000
	movq	%r13, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movq	48(%rsp), %rdi
	movq	40(%rsp), %rax                  # 8-byte Reload
	movq	(%rax,%r15,8), %rbp
	movl	$2048, %edx                     # imm = 0x800
	movq	%rbp, %rsi
	movl	$1, %ecx
	callq	hipMemcpy
	movabsq	$4294967326, %rdi               # imm = 0x10000001E
	movl	$1, %esi
	movq	24(%rsp), %rdx                  # 8-byte Reload
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB4_13
# %bb.12:                               #   in Loop: Header=BB4_11 Depth=2
	movq	48(%rsp), %rdx
	movq	64(%rsp), %rcx
	movq	56(%rsp), %r8
	xorl	%edi, %edi
	movq	144(%rsp), %rsi                 # 8-byte Reload
	movq	136(%rsp), %r9                  # 8-byte Reload
                                        # kill: def $r9d killed $r9d killed $r9
	pushq	8(%rsp)
	.cfi_adjust_cfa_offset 8
	pushq	8(%rsp)
	.cfi_adjust_cfa_offset 8
	callq	_Z32__device_stub__cuda_stringsearchllPcS_PiiPlS_
	addq	$16, %rsp
	.cfi_adjust_cfa_offset -16
.LBB4_13:                               #   in Loop: Header=BB4_11 Depth=2
	callq	hipDeviceSynchronize
	movq	(%rsp), %rsi
	movl	$245760, %edx                   # imm = 0x3C000
	movq	%r13, %rdi
	movl	$2, %ecx
	callq	hipMemcpy
	movq	(%rsp), %rdi
	callq	hipFree
	movq	48(%rsp), %rdi
	callq	hipFree
	xorl	%eax, %eax
	xorl	%ecx, %ecx
.LBB4_14:                               # %.lr.ph.i159
                                        #   Parent Loop BB4_9 Depth=1
                                        #     Parent Loop BB4_11 Depth=2
                                        # =>    This Inner Loop Header: Depth=3
	addq	(%r13,%rax,8), %rcx
	incq	%rax
	cmpq	$30720, %rax                    # imm = 0x7800
	jne	.LBB4_14
# %bb.15:                               # %_Z9sumVectorPll.exit
                                        #   in Loop: Header=BB4_11 Depth=2
	movq	32(%rsp), %rax                  # 8-byte Reload
	addq	%rcx, (%rax,%r15,8)
	movl	$.L.str.12, %edi
	movl	%r15d, %esi
	movq	%rbp, %rdx
	xorl	%eax, %eax
	callq	printf
	movq	%r13, %rdi
	callq	free
	incq	%r15
	cmpq	%r14, %r15
	jne	.LBB4_11
.LBB4_16:                               # %._crit_edge
                                        #   in Loop: Header=BB4_9 Depth=1
	callq	hipPeekAtLastError
	testl	%eax, %eax
	jne	.LBB4_17
# %bb.19:                               # %_Z9gpuAssert10hipError_tPKcib.exit
                                        #   in Loop: Header=BB4_9 Depth=1
	callq	hipDeviceSynchronize
	testl	%eax, %eax
	movq	16(%rsp), %r15                  # 8-byte Reload
	jne	.LBB4_20
# %bb.21:                               # %_Z9gpuAssert10hipError_tPKcib.exit158
                                        #   in Loop: Header=BB4_9 Depth=1
	movq	8(%rsp), %rsi
	movq	%rbx, %rdi
	movq	104(%rsp), %rdx                 # 8-byte Reload
	movl	$2, %ecx
	callq	hipMemcpy
	movq	64(%rsp), %rdi
	callq	hipFree
	movq	56(%rsp), %rdi
	callq	hipFree
	movq	8(%rsp), %rdi
	callq	hipFree
	movl	$.Lstr.2, %edi
	callq	puts@PLT
	testq	%r15, %r15
	movq	80(%rsp), %r13                  # 8-byte Reload
	js	.LBB4_27
# %bb.22:                               # %.lr.ph178.preheader
                                        #   in Loop: Header=BB4_9 Depth=1
	incq	%r12
	xorl	%r14d, %r14d
.LBB4_23:                               # %.lr.ph178
                                        #   Parent Loop BB4_9 Depth=1
                                        # =>  This Inner Loop Header: Depth=2
	movzbl	(%rbx,%r14), %eax
	testl	%eax, %eax
	je	.LBB4_26
# %bb.24:                               # %.lr.ph178
                                        #   in Loop: Header=BB4_23 Depth=2
	cmpl	$36, %eax
	je	.LBB4_26
# %bb.25:                               #   in Loop: Header=BB4_23 Depth=2
	movsbl	%al, %edi
	movq	%r13, %rsi
	callq	fputc@PLT
.LBB4_26:                               #   in Loop: Header=BB4_23 Depth=2
	incq	%r14
	cmpq	%r14, %r12
	jne	.LBB4_23
.LBB4_27:                               # %._crit_edge179
                                        #   in Loop: Header=BB4_9 Depth=1
	movl	$10, %edi
	callq	putchar@PLT
	movl	$.Lstr.3, %edi
	callq	puts@PLT
	leaq	-300000000(%r15), %r14
	movq	%rbx, %rdi
	callq	free
	movl	72(%rsp), %ebp                  # 4-byte Reload
	cmpq	$300000000, %r15                # imm = 0x11E1A300
	movq	%r14, %r15
	movq	96(%rsp), %r13                  # 8-byte Reload
	jg	.LBB4_9
# %bb.28:                               # %.preheader
	movq	%r13, %r14
	movq	88(%rsp), %r13                  # 8-byte Reload
	testl	%r13d, %r13d
	movq	40(%rsp), %r15                  # 8-byte Reload
	movq	32(%rsp), %r12                  # 8-byte Reload
	jle	.LBB4_31
# %bb.29:                               # %.lr.ph182.preheader
	incl	%r13d
	movl	$1, %ebx
.LBB4_30:                               # %.lr.ph182
                                        # =>This Inner Loop Header: Depth=1
	movq	(%r15,%rbx,8), %rdx
	movq	(%r12,%rbx,8), %rcx
	movl	$.L.str.17, %edi
	movl	%ebx, %esi
	xorl	%eax, %eax
	callq	printf
	incq	%rbx
	cmpq	%rbx, %r13
	jne	.LBB4_30
.LBB4_31:                               # %._crit_edge183
	movq	%r14, %rdi
	callq	fclose
	movq	80(%rsp), %rdi                  # 8-byte Reload
	callq	fclose
	xorl	%edi, %edi
	callq	time
	subq	112(%rsp), %rax                 # 8-byte Folded Reload
	cvtsi2sd	%rax, %xmm0
	movl	$.L.str.18, %edi
	movb	$1, %al
	callq	printf
	xorl	%eax, %eax
	addq	$152, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%r12
	.cfi_def_cfa_offset 40
	popq	%r13
	.cfi_def_cfa_offset 32
	popq	%r14
	.cfi_def_cfa_offset 24
	popq	%r15
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	retq
.LBB4_17:
	.cfi_def_cfa_offset 208
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.19, %esi
	movl	$.L.str.13, %ecx
	movq	%rbx, %rdi
	movq	%rax, %rdx
	movl	$356, %r8d                      # imm = 0x164
	jmp	.LBB4_18
.LBB4_20:
	movl	%eax, %ebp
	movq	stderr(%rip), %rbx
	movl	%eax, %edi
	callq	hipGetErrorString
	movl	$.L.str.19, %esi
	movl	$.L.str.13, %ecx
	movq	%rbx, %rdi
	movq	%rax, %rdx
	movl	$357, %r8d                      # imm = 0x165
.LBB4_18:
	xorl	%eax, %eax
	callq	fprintf
	movl	%ebp, %edi
	callq	exit
.LBB4_32:
	movq	stderr(%rip), %rsi
	movl	$.L.str.3, %edi
	callq	fputs
	movl	$1, %edi
	callq	exit
.Lfunc_end4:
	.size	main, .Lfunc_end4-main
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_ctor
	.type	__hip_module_ctor,@function
__hip_module_ctor:                      # @__hip_module_ctor
	.cfi_startproc
# %bb.0:
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	jne	.LBB5_2
# %bb.1:
	movl	$__hip_fatbin_wrapper, %edi
	callq	__hipRegisterFatBinary
	movq	%rax, %rdi
	movq	%rax, __hip_gpubin_handle(%rip)
.LBB5_2:
	xorps	%xmm0, %xmm0
	movups	%xmm0, 16(%rsp)
	movups	%xmm0, (%rsp)
	movl	$_Z17cuda_stringsearchllPcS_PiiPlS_, %esi
	movl	$.L__unnamed_1, %edx
	movl	$.L__unnamed_1, %ecx
	movl	$-1, %r8d
	xorl	%r9d, %r9d
	callq	__hipRegisterFunction
	movl	$__hip_module_dtor, %edi
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	jmp	atexit                          # TAILCALL
.Lfunc_end5:
	.size	__hip_module_ctor, .Lfunc_end5-__hip_module_ctor
	.cfi_endproc
                                        # -- End function
	.p2align	4, 0x90                         # -- Begin function __hip_module_dtor
	.type	__hip_module_dtor,@function
__hip_module_dtor:                      # @__hip_module_dtor
	.cfi_startproc
# %bb.0:
	movq	__hip_gpubin_handle(%rip), %rdi
	testq	%rdi, %rdi
	je	.LBB6_2
# %bb.1:
	pushq	%rax
	.cfi_def_cfa_offset 16
	callq	__hipUnregisterFatBinary
	movq	$0, __hip_gpubin_handle(%rip)
	addq	$8, %rsp
	.cfi_def_cfa_offset 8
.LBB6_2:
	retq
.Lfunc_end6:
	.size	__hip_module_dtor, .Lfunc_end6-__hip_module_dtor
	.cfi_endproc
                                        # -- End function
	.type	WordCount,@object               # @WordCount
	.bss
	.globl	WordCount
	.p2align	2, 0x0
WordCount:
	.long	0                               # 0x0
	.size	WordCount, 4

	.type	_Z17cuda_stringsearchllPcS_PiiPlS_,@object # @_Z17cuda_stringsearchllPcS_PiiPlS_
	.section	.rodata,"a",@progbits
	.globl	_Z17cuda_stringsearchllPcS_PiiPlS_
	.p2align	3, 0x0
_Z17cuda_stringsearchllPcS_PiiPlS_:
	.quad	_Z32__device_stub__cuda_stringsearchllPcS_PiiPlS_
	.size	_Z17cuda_stringsearchllPcS_PiiPlS_, 8

	.type	.L.str,@object                  # @.str
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str:
	.asciz	"%ld "
	.size	.L.str, 5

	.type	_ZL8FileName,@object            # @_ZL8FileName
	.section	.rodata,"a",@progbits
_ZL8FileName:
	.asciz	"01.nt"
	.size	_ZL8FileName, 6

	.type	.L.str.2,@object                # @.str.2
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.2:
	.asciz	"r"
	.size	.L.str.2, 2

	.type	.L.str.3,@object                # @.str.3
.L.str.3:
	.asciz	"File error"
	.size	.L.str.3, 11

	.type	.L.str.4,@object                # @.str.4
.L.str.4:
	.asciz	"Buffer index size %lu \n"
	.size	.L.str.4, 24

	.type	_ZL16QuestionFileName,@object   # @_ZL16QuestionFileName
	.section	.rodata,"a",@progbits
_ZL16QuestionFileName:
	.asciz	"line1"
	.size	_ZL16QuestionFileName, 6

	.type	.L.str.5,@object                # @.str.5
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.5:
	.asciz	"Question Buffer index size %lu \n"
	.size	.L.str.5, 33

	.type	.L.str.7,@object                # @.str.7
.L.str.7:
	.asciz	"Question max length : %d\n"
	.size	.L.str.7, 26

	.type	.L.str.8,@object                # @.str.8
.L.str.8:
	.asciz	"Question elements count : %d\n"
	.size	.L.str.8, 30

	.type	.L.str.9,@object                # @.str.9
.L.str.9:
	.asciz	"Overflow String size : %d\n"
	.size	.L.str.9, 27

	.type	_ZL14OutputFileName,@object     # @_ZL14OutputFileName
	.section	.rodata,"a",@progbits
_ZL14OutputFileName:
	.asciz	"v.txt"
	.size	_ZL14OutputFileName, 6

	.type	.L.str.10,@object               # @.str.10
	.section	.rodata.str1.1,"aMS",@progbits,1
.L.str.10:
	.asciz	"a"
	.size	.L.str.10, 2

	.type	.L.str.11,@object               # @.str.11
.L.str.11:
	.asciz	"This will send buffer start at %ld to %ld of all %ld\n"
	.size	.L.str.11, 54

	.type	.L.str.12,@object               # @.str.12
.L.str.12:
	.asciz	"HOST :: Finish iteration %d at question : %s  temporary founded %ld\n\n"
	.size	.L.str.12, 70

	.type	.L.str.13,@object               # @.str.13
.L.str.13:
	.asciz	"/home/ubuntu/Datasets/stackv2/train-structured-repos-hip-sm89/XWinterVarit/LearnParallel/master/cuda3.hip"
	.size	.L.str.13, 106

	.type	.L.str.17,@object               # @.str.17
.L.str.17:
	.asciz	"element at : %d is : %s finally founded : %lu\n"
	.size	.L.str.17, 47

	.type	.L.str.18,@object               # @.str.18
.L.str.18:
	.asciz	"\nestimate using time : %.2f\n"
	.size	.L.str.18, 29

	.type	.L.str.19,@object               # @.str.19
.L.str.19:
	.asciz	"GPUassert: %s %s %d\n"
	.size	.L.str.19, 21

	.type	.L__unnamed_1,@object           # @0
.L__unnamed_1:
	.asciz	"_Z17cuda_stringsearchllPcS_PiiPlS_"
	.size	.L__unnamed_1, 35

	.type	__hip_fatbin_wrapper,@object    # @__hip_fatbin_wrapper
	.section	.hipFatBinSegment,"a",@progbits
	.p2align	3, 0x0
__hip_fatbin_wrapper:
	.long	1212764230                      # 0x48495046
	.long	1                               # 0x1
	.quad	__hip_fatbin
	.quad	0
	.size	__hip_fatbin_wrapper, 24

	.type	__hip_gpubin_handle,@object     # @__hip_gpubin_handle
	.local	__hip_gpubin_handle
	.comm	__hip_gpubin_handle,8,8
	.section	.init_array,"aw",@init_array
	.p2align	3, 0x0
	.quad	__hip_module_ctor
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.type	.Lstr.1,@object                 # @str.1
	.section	.rodata.str1.1,"aMS",@progbits,1
.Lstr.1:
	.asciz	"This is question file --------"
	.size	.Lstr.1, 31

	.type	.Lstr.2,@object                 # @str.2
.Lstr.2:
	.asciz	"---------saving change buffered----------------------------------------------------------------"
	.size	.Lstr.2, 96

	.type	.Lstr.3,@object                 # @str.3
.Lstr.3:
	.asciz	"-------------------------------------------------------------------------------------------"
	.size	.Lstr.3, 92

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym _Z32__device_stub__cuda_stringsearchllPcS_PiiPlS_
	.addrsig_sym __hip_module_ctor
	.addrsig_sym __hip_module_dtor
	.addrsig_sym _Z17cuda_stringsearchllPcS_PiiPlS_
	.addrsig_sym _ZL8FileName
	.addrsig_sym _ZL16QuestionFileName
	.addrsig_sym _ZL14OutputFileName
	.addrsig_sym __hip_fatbin
	.addrsig_sym __hip_fatbin_wrapper
	.addrsig_sym __hip_cuid_

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.protected	_Z17cuda_stringsearchllPcS_PiiPlS_ ; -- Begin function _Z17cuda_stringsearchllPcS_PiiPlS_
	.globl	_Z17cuda_stringsearchllPcS_PiiPlS_
	.p2align	8
	.type	_Z17cuda_stringsearchllPcS_PiiPlS_,@function
_Z17cuda_stringsearchllPcS_PiiPlS_:     ; @_Z17cuda_stringsearchllPcS_PiiPlS_
; %bb.0:
	s_clause 0x1
	s_load_b32 s2, s[0:1], 0x4c
	s_load_b128 s[4:7], s[0:1], 0x8
	s_waitcnt lgkmcnt(0)
	s_and_b32 s2, s2, 0xffff
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(SKIP_1) | instid1(VALU_DEP_1)
	v_mad_u64_u32 v[1:2], null, s15, s2, v[0:1]
	s_mov_b32 s2, exec_lo
	v_mad_i64_i32 v[3:4], null, 0x2710, v1, 0
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_ge_i64_e64 s[4:5], v[3:4]
	s_cbranch_execz .LBB0_32
; %bb.1:
	s_load_b128 s[8:11], s[0:1], 0x30
	v_add_co_u32 v6, vcc_lo, 0x270f, v3
	v_add_co_ci_u32_e32 v7, vcc_lo, 0, v4, vcc_lo
	v_dual_mov_b32 v5, 0 :: v_dual_and_b32 v4, -16, v3
	v_ashrrev_i32_e32 v2, 31, v1
	s_delay_alu instid0(VALU_DEP_3)
	v_cmp_gt_i64_e32 vcc_lo, s[4:5], v[6:7]
	v_mov_b32_e32 v8, 0
	v_mov_b32_e32 v9, 0
	v_cndmask_b32_e32 v7, s5, v7, vcc_lo
	v_cndmask_b32_e32 v6, s4, v6, vcc_lo
	s_mov_b32 s4, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_le_i64_e64 v[4:5], v[6:7]
	s_cbranch_execz .LBB0_31
; %bb.2:                                ; %.lr.ph52.preheader.i
	global_load_u8 v11, v5, s[6:7]
	s_load_b64 s[0:1], s[0:1], 0x18
	v_mov_b32_e32 v0, 36
	v_mov_b32_e32 v8, v5
	s_add_u32 s5, s6, 1
	s_addc_u32 s13, s7, 0
	s_mov_b32 s12, 0
.LBB0_3:                                ; %.lr.ph52.i
                                        ; =>This Loop Header: Depth=1
                                        ;     Child Loop BB0_5 Depth 2
                                        ;     Child Loop BB0_8 Depth 2
                                        ;     Child Loop BB0_16 Depth 2
                                        ;     Child Loop BB0_23 Depth 2
                                        ;       Child Loop BB0_24 Depth 3
	s_waitcnt vmcnt(0)
	v_dual_mov_b32 v9, 0 :: v_dual_and_b32 v12, 0xff, v11
	v_mov_b32_e32 v10, 0
	s_mov_b32 s14, exec_lo
	s_delay_alu instid0(VALU_DEP_2)
	v_cmpx_ne_u16_e32 0, v12
	s_cbranch_execz .LBB0_7
; %bb.4:                                ; %tailrecurse.i.i.preheader
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_mov_b64 s[2:3], 0
.LBB0_5:                                ; %tailrecurse.i.i
                                        ;   Parent Loop BB0_3 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_delay_alu instid0(SALU_CYCLE_1)
	s_add_u32 s16, s5, s2
	s_addc_u32 s17, s13, s3
	s_add_u32 s2, s2, 1
	global_load_u8 v9, v5, s[16:17]
	s_addc_u32 s3, s3, 0
	s_waitcnt vmcnt(0)
	v_cmp_ne_u16_e32 vcc_lo, 0, v9
	s_cbranch_vccnz .LBB0_5
; %bb.6:                                ; %Flow78
                                        ;   in Loop: Header=BB0_3 Depth=1
	v_dual_mov_b32 v10, s3 :: v_dual_mov_b32 v9, s2
.LBB0_7:                                ; %Flow79
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s14
	s_waitcnt lgkmcnt(0)
	v_add_co_u32 v12, vcc_lo, s0, v4
	v_add_co_ci_u32_e32 v13, vcc_lo, s1, v5, vcc_lo
	s_mov_b64 s[2:3], 0
	s_mov_b32 s15, 0
                                        ; implicit-def: $sgpr14
                                        ; implicit-def: $sgpr17
                                        ; implicit-def: $sgpr16
.LBB0_8:                                ; %_Z8d_strlenPKc.exit.i
                                        ;   Parent Loop BB0_3 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	s_or_b32 s16, s16, exec_lo
	s_or_b32 s17, s17, exec_lo
	s_mov_b32 s18, exec_lo
                                        ; implicit-def: $vgpr14
                                        ; implicit-def: $vgpr15
	v_cmpx_ne_u64_e64 s[2:3], v[9:10]
	s_cbranch_execz .LBB0_10
; %bb.9:                                ;   in Loop: Header=BB0_8 Depth=2
	v_add_co_u32 v14, vcc_lo, v12, s2
	v_add_co_ci_u32_e32 v15, vcc_lo, s3, v13, vcc_lo
	s_add_u32 s20, s6, s2
	s_addc_u32 s21, s7, s3
	s_add_u32 s2, s2, 1
	global_load_u8 v14, v[14:15], off
	global_load_u8 v15, v5, s[20:21]
	s_addc_u32 s3, s3, 0
	s_and_not1_b32 s17, s17, exec_lo
	s_and_not1_b32 s16, s16, exec_lo
	s_waitcnt vmcnt(0)
	v_cmp_ne_u16_e32 vcc_lo, v14, v15
	s_and_b32 s19, vcc_lo, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s17, s17, s19
.LBB0_10:                               ; %Flow77
                                        ;   in Loop: Header=BB0_8 Depth=2
	s_or_b32 exec_lo, exec_lo, s18
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s18, exec_lo, s17
	s_or_b32 s15, s18, s15
	s_and_not1_b32 s14, s14, exec_lo
	s_and_b32 s18, s16, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s14, s14, s18
	s_and_not1_b32 exec_lo, exec_lo, s15
	s_cbranch_execnz .LBB0_8
; %bb.11:                               ; %loop.exit.guard
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s15
	v_mov_b32_e32 v9, 0
	s_xor_b32 s2, s14, -1
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_saveexec_b32 s3, s2
	s_xor_b32 s2, exec_lo, s3
; %bb.12:                               ;   in Loop: Header=BB0_3 Depth=1
	v_and_b32_e32 v9, 0xff, v14
	v_and_b32_e32 v10, 0xff, v15
	s_delay_alu instid0(VALU_DEP_1)
	v_sub_nc_u32_e32 v9, v9, v10
; %bb.13:                               ; %_Z9d_strncmpPKcS0_m.exit.i
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	s_mov_b32 s2, exec_lo
	v_cmpx_ne_u32_e32 0, v9
	s_xor_b32 s2, exec_lo, s2
	s_cbranch_execz .LBB0_21
; %bb.14:                               ; %.preheader.i
                                        ;   in Loop: Header=BB0_3 Depth=1
	v_ashrrev_i32_e32 v4, 31, v3
	s_mov_b32 s3, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_ge_i64_e64 v[6:7], v[3:4]
	s_cbranch_execz .LBB0_20
; %bb.15:                               ; %.lr.ph.i.preheader
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_mov_b32 s14, 0
                                        ; implicit-def: $sgpr15
.LBB0_16:                               ; %.lr.ph.i
                                        ;   Parent Loop BB0_3 Depth=1
                                        ; =>  This Inner Loop Header: Depth=2
	v_add_co_u32 v9, vcc_lo, s0, v3
	v_add_co_ci_u32_e32 v10, vcc_lo, s1, v4, vcc_lo
	s_or_b32 s15, s15, exec_lo
	s_mov_b32 s16, exec_lo
	global_load_u8 v4, v[9:10], off
	s_waitcnt vmcnt(0)
	v_cmpx_ne_u16_e32 10, v4
; %bb.17:                               ;   in Loop: Header=BB0_16 Depth=2
	v_add_nc_u32_e32 v3, 1, v3
	s_and_not1_b32 s15, s15, exec_lo
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_1)
	v_ashrrev_i32_e32 v4, 31, v3
	v_cmp_lt_i64_e32 vcc_lo, v[6:7], v[3:4]
	s_and_b32 s17, vcc_lo, exec_lo
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 s15, s15, s17
; %bb.18:                               ; %Flow
                                        ;   in Loop: Header=BB0_16 Depth=2
	s_or_b32 exec_lo, exec_lo, s16
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s16, exec_lo, s15
	s_or_b32 s14, s16, s14
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s14
	s_cbranch_execnz .LBB0_16
; %bb.19:                               ; %Flow70
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s14
.LBB0_20:                               ; %Flow71
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_delay_alu instid0(SALU_CYCLE_1)
	s_or_b32 exec_lo, exec_lo, s3
.LBB0_21:                               ; %Flow75
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_and_not1_saveexec_b32 s14, s2
	s_cbranch_execz .LBB0_29
; %bb.22:                               ; %.preheader.preheader
                                        ;   in Loop: Header=BB0_3 Depth=1
	v_mov_b32_e32 v9, v3
	s_mov_b32 s15, 0
.LBB0_23:                               ; %.preheader
                                        ;   Parent Loop BB0_3 Depth=1
                                        ; =>  This Loop Header: Depth=2
                                        ;       Child Loop BB0_24 Depth 3
	global_load_u8 v11, v5, s[6:7]
	s_mov_b64 s[2:3], 0
	s_waitcnt vmcnt(0)
	v_cmp_eq_u16_e32 vcc_lo, 0, v11
	s_cbranch_vccnz .LBB0_25
.LBB0_24:                               ; %tailrecurse.i35.i
                                        ;   Parent Loop BB0_3 Depth=1
                                        ;     Parent Loop BB0_23 Depth=2
                                        ; =>    This Inner Loop Header: Depth=3
	s_add_u32 s16, s5, s2
	s_addc_u32 s17, s13, s3
	s_add_u32 s2, s2, 1
	global_load_u8 v10, v5, s[16:17]
	s_addc_u32 s3, s3, 0
	s_waitcnt vmcnt(0)
	v_cmp_ne_u16_e32 vcc_lo, 0, v10
	s_cbranch_vccnz .LBB0_24
.LBB0_25:                               ; %_Z8d_strlenPKc.exit40.i
                                        ;   in Loop: Header=BB0_23 Depth=2
	v_add_co_u32 v12, vcc_lo, s2, v4
	v_ashrrev_i32_e32 v10, 31, v9
	v_add_co_ci_u32_e32 v13, vcc_lo, s3, v5, vcc_lo
	s_mov_b32 s3, -1
	s_mov_b32 s2, exec_lo
	s_delay_alu instid0(VALU_DEP_1)
	v_cmpx_gt_u64_e64 v[12:13], v[9:10]
	s_cbranch_execz .LBB0_27
; %bb.26:                               ;   in Loop: Header=BB0_23 Depth=2
	v_add_co_u32 v12, vcc_lo, s10, v9
	v_add_co_ci_u32_e32 v13, vcc_lo, s11, v10, vcc_lo
	v_add_nc_u32_e32 v9, 1, v9
	s_xor_b32 s3, exec_lo, -1
	global_store_b8 v[12:13], v0, off
.LBB0_27:                               ; %Flow74
                                        ;   in Loop: Header=BB0_23 Depth=2
	s_or_b32 exec_lo, exec_lo, s2
	s_delay_alu instid0(SALU_CYCLE_1) | instskip(NEXT) | instid1(SALU_CYCLE_1)
	s_and_b32 s2, exec_lo, s3
	s_or_b32 s15, s2, s15
	s_delay_alu instid0(SALU_CYCLE_1)
	s_and_not1_b32 exec_lo, exec_lo, s15
	s_cbranch_execnz .LBB0_23
; %bb.28:                               ; %.loopexit.loopexit.i
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s15
	v_add_nc_u32_e32 v8, 1, v8
.LBB0_29:                               ; %Flow76
                                        ;   in Loop: Header=BB0_3 Depth=1
	s_or_b32 exec_lo, exec_lo, s14
	v_add_nc_u32_e32 v4, 1, v3
	s_delay_alu instid0(VALU_DEP_1) | instskip(SKIP_2) | instid1(SALU_CYCLE_1)
	v_cmp_lt_i64_e32 vcc_lo, v[6:7], v[4:5]
	v_mov_b32_e32 v3, v4
	s_or_b32 s12, vcc_lo, s12
	s_and_not1_b32 exec_lo, exec_lo, s12
	s_cbranch_execnz .LBB0_3
; %bb.30:                               ; %_Z16string_search_rrllPcS_icS_.exit.loopexit
	s_or_b32 exec_lo, exec_lo, s12
	v_ashrrev_i32_e32 v9, 31, v8
.LBB0_31:                               ; %Flow80
	s_or_b32 exec_lo, exec_lo, s4
	v_lshlrev_b64 v[0:1], 3, v[1:2]
	s_waitcnt lgkmcnt(0)
	s_delay_alu instid0(VALU_DEP_1) | instskip(NEXT) | instid1(VALU_DEP_2)
	v_add_co_u32 v0, vcc_lo, s8, v0
	v_add_co_ci_u32_e32 v1, vcc_lo, s9, v1, vcc_lo
	global_store_b64 v[0:1], v[8:9], off
.LBB0_32:
	s_nop 0
	s_sendmsg sendmsg(MSG_DEALLOC_VGPRS)
	s_endpgm
	.section	.rodata,"a",@progbits
	.p2align	6, 0x0
	.amdhsa_kernel _Z17cuda_stringsearchllPcS_PiiPlS_
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 320
		.amdhsa_user_sgpr_count 15
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_wavefront_size32 1
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_enable_private_segment 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 16
		.amdhsa_next_free_sgpr 22
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_workgroup_processor_mode 1
		.amdhsa_memory_ordered 1
		.amdhsa_forward_progress 0
		.amdhsa_shared_vgpr_count 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z17cuda_stringsearchllPcS_PiiPlS_, .Lfunc_end0-_Z17cuda_stringsearchllPcS_PiiPlS_
                                        ; -- End function
	.section	.AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 924
; NumSgprs: 24
; NumVgprs: 16
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 2
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 24
; NumVGPRsForWavesPerEU: 16
; Occupancy: 16
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 15
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .offset:         0
        .size:           8
        .value_kind:     by_value
      - .offset:         8
        .size:           8
        .value_kind:     by_value
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         24
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         32
        .size:           8
        .value_kind:     global_buffer
      - .offset:         40
        .size:           4
        .value_kind:     by_value
      - .address_space:  global
        .offset:         48
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         56
        .size:           8
        .value_kind:     global_buffer
      - .offset:         64
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         68
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         72
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         76
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         78
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         80
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         82
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         84
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         86
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         104
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         112
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         120
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         128
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 320
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z17cuda_stringsearchllPcS_PiiPlS_
    .private_segment_fixed_size: 0
    .sgpr_count:     24
    .sgpr_spill_count: 0
    .symbol:         _Z17cuda_stringsearchllPcS_PiiPlS_.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     16
    .vgpr_spill_count: 0
    .wavefront_size: 32
    .workgroup_processor_mode: 1
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
