Ejemplo n.º 1
0
CUresult
TestSAXPY( chCUDADevice *chDevice, size_t N, float alpha )
{
    CUresult status;
    CUdeviceptr dptrOut = 0;
    CUdeviceptr dptrIn = 0;
    float *hostOut = 0;
    float *hostIn = 0;

    CUDA_CHECK( cuCtxPushCurrent( chDevice->context() ) );

    CUDA_CHECK( cuMemAlloc( &dptrOut, N*sizeof(float) ) );
    CUDA_CHECK( cuMemsetD32( dptrOut, 0, N ) );
    CUDA_CHECK( cuMemAlloc( &dptrIn, N*sizeof(float) ) );
    CUDA_CHECK( cuMemHostAlloc( (void **) &hostOut, N*sizeof(float), 0 ) );
    CUDA_CHECK( cuMemHostAlloc( (void **) &hostIn, N*sizeof(float), 0 ) );
    for ( size_t i = 0; i < N; i++ ) {
        hostIn[i] = (float) rand() / (float) RAND_MAX;
    }
    CUDA_CHECK( cuMemcpyHtoDAsync( dptrIn, hostIn, N*sizeof(float ), NULL ) );

    {
        CUmodule moduleSAXPY;
        CUfunction kernelSAXPY;
        void *params[] = { &dptrOut, &dptrIn, &N, &alpha };
        
        moduleSAXPY = chDevice->module( "saxpy.ptx" );
        if ( ! moduleSAXPY ) {
            status = CUDA_ERROR_NOT_FOUND;
            goto Error;
        }
        CUDA_CHECK( cuModuleGetFunction( &kernelSAXPY, moduleSAXPY, "saxpy" ) );

        CUDA_CHECK( cuLaunchKernel( kernelSAXPY, 1500, 1, 1, 512, 1, 1, 0, NULL, params, NULL ) );

    }

    CUDA_CHECK( cuMemcpyDtoHAsync( hostOut, dptrOut, N*sizeof(float), NULL ) );
    CUDA_CHECK( cuCtxSynchronize() );
    for ( size_t i = 0; i < N; i++ ) {
        if ( fabsf( hostOut[i] - alpha*hostIn[i] ) > 1e-5f ) {
            status = CUDA_ERROR_UNKNOWN;
            goto Error;
        }
    }
    status = CUDA_SUCCESS;
    printf( "Well it worked!\n" );

Error:
    cuCtxPopCurrent( NULL );
    cuMemFreeHost( hostOut );
    cuMemFreeHost( hostIn );
    cuMemFree( dptrOut );
    cuMemFree( dptrIn );
    return status;
}
Ejemplo n.º 2
0
void* GPUInterface::AllocatePinnedHostMemory(size_t memSize, bool writeCombined, bool mapped) {
#ifdef BEAGLE_DEBUG_FLOW
    fprintf(stderr,"\t\t\tEntering GPUInterface::AllocatePinnedHostMemory\n");
#endif

    void* ptr;

    unsigned int flags = 0;

    if (writeCombined)
        flags |= CU_MEMHOSTALLOC_WRITECOMBINED;
    if (mapped)
        flags |= CU_MEMHOSTALLOC_DEVICEMAP;

    SAFE_CUPP(cuMemHostAlloc(&ptr, memSize, flags));


#ifdef BEAGLE_DEBUG_VALUES
    fprintf(stderr, "Allocated pinned host (CPU) memory %ld to %lu .\n", (long)ptr, ((long)ptr + memSize));
#endif

#ifdef BEAGLE_DEBUG_FLOW
    fprintf(stderr, "\t\t\tLeaving  GPUInterface::AllocatePinnedHostMemory\n");
#endif

    return ptr;
}
int gib_alloc ( void **buffers, int buf_size, int *ld, gib_context c ) {
  ERROR_CHECK_FAIL(cuCtxPushCurrent(((gpu_context)(c->acc_context))->pCtx));
#if GIB_USE_MMAP
  ERROR_CHECK_FAIL(cuMemHostAlloc(buffers, (c->n+c->m)*buf_size, 
				  CU_MEMHOSTALLOC_DEVICEMAP));
#else
  ERROR_CHECK_FAIL(cuMemAllocHost(buffers, (c->n+c->m)*buf_size));
#endif
  *ld = buf_size;
  ERROR_CHECK_FAIL(cuCtxPopCurrent(&((gpu_context)(c->acc_context))->pCtx));
  return GIB_SUC;
}
Ejemplo n.º 4
0
int main(int argc, char *argv[])
{
	char c;
	CUcontext ctx;
	CUdevice dev = 0;
	void *toSpace;
	int status, free, total;
	CUdeviceptr ptr = (CUdeviceptr)NULL;
	int size;
	
	if(argc != 2){
		fprintf(stderr,"Usage: mem_alloc.exe [MEMORY TO ALLOCATE IN MB]\n");
		exit(1);
	}
	
	printf("All status results should be 0, if not an error has occured.\nIf 2 is reported an out of memory error has occured for\nwhich you should decrease the memory input\n");
	size = atoi(argv[1]);
	
	printf("\nTrying to allocate %iMB of memory on host and GPU\n",size);
	
	if(size <= 0){
		fprintf(stderr,"\nERROR: Memory must be greater than 0\n");
		exit(1);
	}
	
	status = cuInit(0);
	printf("Init status: %i\n",status); 

	status = cuCtxCreate(&ctx, 0, dev);
	printf("Context creation status: %i\n",status); 
	
	cuMemGetInfo(&free, &total);
	printf("Get memory info status: %i\n",status); 
	
	printf("\n%.1f/%.1f (Free/Total) MB\n", free/1024.0/1024.0, total/1024.0/1024.0);
	
	status = cuMemHostAlloc(&toSpace, size*1024*1024, 0); 
	printf("Host allocation status: %i %s\n",status, (status==CUDA_SUCCESS) ? "SUCCESS" : "FAILED"); 

	status = cuMemAlloc(&ptr, size*1024*1024);
	printf("GPU allocation status: %i %s\n",status, (status==CUDA_SUCCESS) ? "SUCCESS" : "FAILED");

	printf("\nPress any key to exit...");
	scanf("%c", &c);
	
	status = cuCtxDestroy(ctx);
	printf("Context destroy status: %i\n",status); 

	return 0;
}
Ejemplo n.º 5
0
SEXP
R_auto_cuMemHostAlloc(SEXP r_bytesize, SEXP r_Flags)
{
    SEXP r_ans = R_NilValue;
    void * pp;
    size_t bytesize = REAL(r_bytesize)[0];
    unsigned int Flags = REAL(r_Flags)[0];
    CUresult ans;
    ans = cuMemHostAlloc(& pp,  bytesize,  Flags);
    if(ans)
       return(R_cudaErrorInfo(ans));
    r_ans = R_createRef(pp, "voidPtr") ;
    return(r_ans);
}
Ejemplo n.º 6
0
void *swanMallocHost( size_t len ) {
	CUresult err;
	void *ptr;
	try_init();

#ifdef DEV_EXTENSIONS
	err= cuMemHostAlloc( &ptr, len, CU_MEMHOSTALLOC_PORTABLE  ); //| CU_MEMHOSTALLOC_DEVICEMAP  ); //| CU_MEMHOSTALLOC_WRITECOMBINED );
#else
	err = cuMemAllocHost( &ptr, len );
#endif

	if ( err != CUDA_SUCCESS ) {
		fprintf( stderr, "swanMallocHost error: %d\n", err );
		error("swanMallocHost failed\n" );
	}
//	printf("MallocHost %p\n", ptr );
	memset( ptr, 0, len );
	return ptr;
}
Ejemplo n.º 7
0
static void calc_a_score_GPU(FLOAT *ac_score,  FLOAT **score,
			     int *ssize_start,  Model_info *MI,
			     FLOAT scale, int *size_score_array,
			     int NoC)
{
	CUresult res;

	const int IHEI = MI->IM_HEIGHT;
	const int IWID = MI->IM_WIDTH;
	int pady_n = MI->pady;
	int padx_n = MI->padx;
	int block_pad = (int)(scale/2.0);

	struct timeval tv;

	int *RY_array, *RX_array;
	res = cuMemHostAlloc((void**)&RY_array, NoC*sizeof(int), CU_MEMHOSTALLOC_DEVICEMAP);
	if(res != CUDA_SUCCESS) {
		printf("cuMemHostAlloc(RY_array) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuMemHostAlloc((void**)&RX_array, NoC*sizeof(int), CU_MEMHOSTALLOC_DEVICEMAP);
	if(res != CUDA_SUCCESS) {
		printf("cuMemHostAlloc(RX_array) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	for(int i = 0; i < NoC; i++) {
		int rsize[2] = {MI->rsize[i*2], MI->rsize[i*2+1]};

		RY_array[i] = (int)((FLOAT)rsize[0]*scale/2.0-1.0+block_pad);
		RX_array[i] = (int)((FLOAT)rsize[1]*scale/2.0-1.0+block_pad);
	}

	CUdeviceptr ac_score_dev, score_dev;
	CUdeviceptr ssize_dev, size_score_dev;
	CUdeviceptr RY_dev, RX_dev;

	int size_score=0;
	for(int i = 0; i < NoC; i++) {
		size_score += size_score_array[i];
	}

	/* allocate GPU memory */
	res = cuMemAlloc(&ac_score_dev, gpu_size_A_SCORE);
	if(res != CUDA_SUCCESS) {
		printf("cuMemAlloc(ac_score) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuMemAlloc(&score_dev, size_score);
	if(res != CUDA_SUCCESS) {
		printf("cuMemAlloc(score) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuMemAlloc(&ssize_dev, NoC*sizeof(int));
	if(res != CUDA_SUCCESS) {
		printf("cuMemAlloc(ssize) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuMemAlloc(&size_score_dev, NoC*sizeof(int));
	if(res != CUDA_SUCCESS) {
		printf("cuMemAlloc(size_score) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuMemAlloc(&RY_dev, NoC*sizeof(int));
	if(res != CUDA_SUCCESS) {
		printf("cuMemAlloc(RY) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuMemAlloc(&RX_dev, NoC*sizeof(int));
	if(res != CUDA_SUCCESS) {
		printf("cuMemAlloc(RX) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	gettimeofday(&tv_memcpy_start, nullptr);
	/* upload date to GPU */
	res = cuMemcpyHtoD(ac_score_dev, &ac_score[0], gpu_size_A_SCORE);
	if(res != CUDA_SUCCESS) {
		printf("cuMemcpyHtoD(ac_score) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuMemcpyHtoD(score_dev, &score[0][0], size_score);
	if(res != CUDA_SUCCESS) {
		printf("cuMemcpyHtoD(score) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuMemcpyHtoD(ssize_dev, &ssize_start[0], NoC*sizeof(int));
	if(res != CUDA_SUCCESS) {
		printf("cuMemcpyHtoD(ssize) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuMemcpyHtoD(size_score_dev, &size_score_array[0], NoC*sizeof(int));
	if(res != CUDA_SUCCESS) {
		printf("cuMemcpyHtoD(size_score) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuMemcpyHtoD(RY_dev, &RY_array[0], NoC*sizeof(int));
	if(res != CUDA_SUCCESS) {
		printf("cuMemcpyHtoD(RY) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuMemcpyHtoD(RX_dev, &RX_array[0], NoC*sizeof(int));
	if(res != CUDA_SUCCESS) {
		printf("cuMemcpyHtoD(RX) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	gettimeofday(&tv_memcpy_end, nullptr);
	tvsub(&tv_memcpy_end, &tv_memcpy_start, &tv);
	time_memcpy += tv.tv_sec * 1000.0 + (float)tv.tv_usec / 1000.0;

	void* kernel_args[] = {
		(void*)&IWID,
		(void*)&IHEI,
		(void*)&scale,
		(void*)&padx_n,
		(void*)&pady_n,
		&RX_dev,
		&RY_dev,
		&ac_score_dev,
		&score_dev,
		&ssize_dev,
		(void*)&NoC,
		&size_score_dev
	};

	int sharedMemBytes = 0;

	/* define CUDA block shape */
	int max_threads_num = 0;
	int thread_num_x, thread_num_y;
	int block_num_x, block_num_y;

	res = cuDeviceGetAttribute(&max_threads_num, CU_DEVICE_ATTRIBUTE_MAX_THREADS_PER_BLOCK, dev[0]);
	if(res != CUDA_SUCCESS){
		printf("\ncuDeviceGetAttribute() failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	NR_MAXTHREADS_X[0] = (int)sqrt((double)max_threads_num/NoC);
	NR_MAXTHREADS_Y[0] = (int)sqrt((double)max_threads_num/NoC);

	thread_num_x = (IWID < NR_MAXTHREADS_X[0]) ? IWID : NR_MAXTHREADS_X[0];
	thread_num_y = (IHEI < NR_MAXTHREADS_Y[0]) ? IHEI : NR_MAXTHREADS_Y[0];

	block_num_x = IWID / thread_num_x;
	block_num_y = IHEI / thread_num_y;
	if(IWID % thread_num_x != 0) block_num_x++;
	if(IHEI % thread_num_y != 0) block_num_y++;

	gettimeofday(&tv_kernel_start, nullptr);
	/* launch GPU kernel */
	res = cuLaunchKernel(
		func_calc_a_score[0], // call function
		block_num_x,       // gridDimX
		block_num_y,       // gridDimY
		1,                 // gridDimZ
		thread_num_x,      // blockDimX
		thread_num_y,      // blockDimY
		NoC,               // blockDimZ
		sharedMemBytes,    // sharedMemBytes
		nullptr,              // hStream
		kernel_args,       // kernelParams
		nullptr               // extra
		);
	if(res != CUDA_SUCCESS) {
		printf("cuLaunchKernel(calc_a_score) failed : res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuCtxSynchronize();
	if(res != CUDA_SUCCESS) {
		printf("cuCtxSynchronize(calc_a_score) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}
	gettimeofday(&tv_kernel_end, nullptr);
	tvsub(&tv_kernel_end, &tv_kernel_start, &tv);
	time_kernel += tv.tv_sec * 1000.0 + (float)tv.tv_usec / 1000.0;

	gettimeofday(&tv_memcpy_start, nullptr);
	/* download data from GPU */
	res = cuMemcpyDtoH(ac_score, ac_score_dev, gpu_size_A_SCORE);
	if(res != CUDA_SUCCESS) {
		printf("cuMemcpyDtoH(ac_score) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	gettimeofday(&tv_memcpy_end, nullptr);
	tvsub(&tv_memcpy_end, &tv_memcpy_start, &tv);
	time_memcpy += tv.tv_sec * 1000.0 + (float)tv.tv_usec / 1000.0;

	/* free GPU memory */
	res = cuMemFree(ac_score_dev);
	if(res != CUDA_SUCCESS) {
		printf("cuMemFree(ac_score_dev) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuMemFree(score_dev);
	if(res != CUDA_SUCCESS) {
		printf("cuMemFree(score_dev) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuMemFree(ssize_dev);
	if(res != CUDA_SUCCESS) {
		printf("cuMemFree(ssize_dev) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuMemFree(size_score_dev);
	if(res != CUDA_SUCCESS) {
		printf("cuMemFree(size_score_dev) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuMemFree(RY_dev);
	if(res != CUDA_SUCCESS) {
		printf("cuMemFree(RY_dev) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuMemFree(RX_dev);
	if(res != CUDA_SUCCESS) {
		printf("cuMemFree(RX_dev) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	/* free CPU memory */
	res = cuMemFreeHost(RY_array);
	if(res != CUDA_SUCCESS) {
		printf("cuMemFreeHost(RY_array) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	res = cuMemFreeHost(RX_array);
	if(res != CUDA_SUCCESS) {
		printf("cuMemFreeHost(RX_array) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}
}
Ejemplo n.º 8
0
//detect boundary box
FLOAT *dpm_ttic_gpu_get_boxes(FLOAT **features,FLOAT *scales,int *feature_size, GPUModel *MO,
			      int *detected_count, FLOAT *acc_score, FLOAT thresh)
{
	//constant parameters
	const int max_scale = MO->MI->max_scale;
	const int interval = MO->MI->interval;
	const int sbin = MO->MI->sbin;
	const int padx = MO->MI->padx;
	const int pady = MO->MI->pady;
	const int NoR = MO->RF->NoR;
	const int NoP = MO->PF->NoP;
	const int NoC = MO->MI->numcomponent;
	const int *numpart = MO->MI->numpart;
	const int LofFeat=(max_scale+interval)*NoC;
	const int L_MAX = max_scale+interval;

	/* for measurement */
	struct timeval tv;
	struct timeval tv_make_c_start, tv_make_c_end;
	struct timeval tv_nucom_start, tv_nucom_end;
	struct timeval tv_box_start, tv_box_end;
	float time_box=0;
	struct timeval tv_root_score_start, tv_root_score_end;
	float time_root_score = 0;
	struct timeval tv_part_score_start, tv_part_score_end;
	float time_part_score = 0;
	struct timeval tv_dt_start, tv_dt_end;
	float time_dt = 0;
	struct timeval tv_calc_a_score_start, tv_calc_a_score_end;
	float time_calc_a_score = 0;

	gettimeofday(&tv_make_c_start, nullptr);

	int **RF_size = MO->RF->root_size;
	int *rootsym = MO->RF->rootsym;
	int *part_sym = MO->PF->part_sym;
	int **part_size = MO->PF->part_size;
	FLOAT **rootfilter = MO->RF->rootfilter;
	FLOAT **partfilter=MO->PF->partfilter;
	int **psize = MO->MI->psize;

	int **rm_size_array = (int **)malloc(sizeof(int *)*L_MAX);
	int **pm_size_array = (int **)malloc(sizeof(int *)*L_MAX);
	pm_size_array = (int **)malloc(sizeof(int *)*L_MAX);

	FLOAT **Tboxes=(FLOAT**)calloc(LofFeat,sizeof(FLOAT*)); //box coordinate information(Temp)
	int  *b_nums =(int*)calloc(LofFeat,sizeof(int)); //length of Tboxes
	int count = 0;
	int detected_boxes=0;
	CUresult res;

	/* matched score (root and part) */
	FLOAT ***rootmatch,***partmatch = nullptr;

	int *new_PADsize;  // need new_PADsize[L_MAX*3]
	size_t SUM_SIZE_feat = 0;

	FLOAT **featp2 = (FLOAT **)malloc(L_MAX*sizeof(FLOAT *));


	if(featp2 == nullptr) {  // error semantics
		printf("allocate featp2 failed\n");
		exit(1);
	}


	/* allocate required memory for new_PADsize */
	new_PADsize = (int *)malloc(L_MAX*3*sizeof(int));
	if(new_PADsize == nullptr) {     // error semantics
		printf("allocate new_PADsize failed\n");
		exit(1);
	}

	/* do padarray once and reuse it at calculating root and part time */

	/* calculate sum of size of padded feature */
	for(int tmpL=0; tmpL<L_MAX; tmpL++) {
		int PADsize[3] = { feature_size[tmpL*2], feature_size[tmpL*2+1], 31 };
		int NEW_Y = PADsize[0] + pady*2;
		int NEW_X = PADsize[1] + padx*2;
		SUM_SIZE_feat += (NEW_X*NEW_Y*PADsize[2])*sizeof(FLOAT);
	}

	/* allocate region for padded feat in a lump */
	FLOAT *dst_feat;
	res = cuMemHostAlloc((void **)&dst_feat, SUM_SIZE_feat, CU_MEMHOSTALLOC_DEVICEMAP);
	if(res != CUDA_SUCCESS) {
		printf("cuMemHostAlloc(dst_feat) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	memset(dst_feat, 0, SUM_SIZE_feat);  // zero clear

	/* distribute allocated region */
	uintptr_t pointer_feat = (uintptr_t)dst_feat;
	for(int tmpL=0; tmpL<L_MAX; tmpL++) {

		featp2[tmpL] = (FLOAT *)pointer_feat;
		int PADsize[3] = { feature_size[tmpL*2], feature_size[tmpL*2+1], 31 };
		int NEW_Y = PADsize[0] + pady*2;
		int NEW_X = PADsize[1] + padx*2;
		pointer_feat += (uintptr_t)(NEW_X*NEW_Y*PADsize[2]*sizeof(FLOAT));

	}

	/* copy feat to feat2 */
	for(int tmpL=0; tmpL<L_MAX; tmpL++) {

		int PADsize[3] = { feature_size[tmpL*2], feature_size[tmpL*2+1], 31 };
		int NEW_Y = PADsize[0] + pady*2;
		int NEW_X = PADsize[1] + padx*2;
		int L = NEW_Y*padx;
		int SPL = PADsize[0] + pady;
		int M_S = sizeof(FLOAT)*PADsize[0];
		FLOAT *P = featp2[tmpL];
		FLOAT *S = features[tmpL];

		for(int i=0; i<PADsize[2]; i++)
		{
			P += L;
			for(int j=0; j<PADsize[1]; j++)
			{
				P += pady;
				memcpy(P, S, M_S);
				S += PADsize[0];
				P += SPL;
			}
			P += L;
		}

		new_PADsize[tmpL*3] = NEW_Y;
		new_PADsize[tmpL*3 + 1] = NEW_X;
		new_PADsize[tmpL*3 + 2] = PADsize[2];

	}

	/* do padarray once and reuse it at calculating root and part time */

	/* allocation in a lump */
	int *dst_rm_size = (int *)malloc(sizeof(int)*NoC*2*L_MAX);
	if(dst_rm_size == nullptr) {
		printf("allocate dst_rm_size failed\n");
		exit(1);
	}

	/* distribution to rm_size_array[L_MAX] */
	uintptr_t ptr = (uintptr_t)dst_rm_size;
	for(int i=0; i<L_MAX; i++) {
		rm_size_array[i] = (int *)ptr;
		ptr += (uintptr_t)(NoC*2*sizeof(int));
	}

	/* allocation in a lump */
	int *dst_pm_size = (int *)malloc(sizeof(int)*NoP*2*L_MAX);
	if(dst_pm_size == nullptr) {
		printf("allocate dst_pm_size failed\n");
		exit(1);
	}

	/* distribution to pm_size_array[L_MAX] */
	ptr = (uintptr_t)dst_pm_size;
	for(int i=0; i<L_MAX; i++) {
		pm_size_array[i] = (int *)ptr;
		ptr += (uintptr_t)(NoP*2*sizeof(int));
	}


	///////level
	for (int level=interval; level<L_MAX; level++)  // feature's loop(A's loop) 1level 1picture
	{
		if(feature_size[level*2]+2*pady<MO->MI->max_Y ||(feature_size[level*2+1]+2*padx<MO->MI->max_X))
		{
			Tboxes[count]=nullptr;
			count++;
			continue;
		}
	}  //for (level)  // feature's loop(A's loop) 1level 1picture

	///////root calculation/////////
	/* calculate model score (only root) */

	gettimeofday(&tv_root_score_start, nullptr);
	rootmatch = fconvsMT_GPU(
		featp2,
		SUM_SIZE_feat,
		rootfilter,
		rootsym,
		1,
		NoR,
		new_PADsize,
		RF_size, rm_size_array,
		L_MAX,
		interval,
		feature_size,
		padx,
		pady,
		MO->MI->max_X,
		MO->MI->max_Y,
		ROOT
		);
	gettimeofday(&tv_root_score_end, nullptr);
	tvsub(&tv_root_score_end, &tv_root_score_start, &tv);
	time_root_score += tv.tv_sec * 1000.0 + (float)tv.tv_usec / 1000.0;

	///////part calculation/////////
	if(NoP>0)
	{
		/* calculate model score (only part) */
		gettimeofday(&tv_part_score_start, nullptr);
		partmatch = fconvsMT_GPU(
			featp2,
			SUM_SIZE_feat,
			partfilter,
			part_sym,
			1,
			NoP,
			new_PADsize,
			part_size,
			pm_size_array,
			L_MAX,
			interval,
			feature_size,
			padx,
			pady,
			MO->MI->max_X,
			MO->MI->max_Y,
			PART
			);
		gettimeofday(&tv_part_score_end, nullptr);
		tvsub(&tv_part_score_end, &tv_part_score_start, &tv);
		time_part_score += tv.tv_sec * 1000.0 + (float)tv.tv_usec / 1000.0;

	}

	res = cuCtxSetCurrent(ctx[0]);
	if(res != CUDA_SUCCESS) {
		printf("cuCtxSetCurrent(ctx[0]) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	gettimeofday(&tv_make_c_end, nullptr);
	gettimeofday(&tv_nucom_start, nullptr);

	count = 0;
	detected_boxes = 0;

	int **RL_array = (int **)malloc((L_MAX-interval)*sizeof(int*));
	int *dst_RL = (int *) malloc(NoC*(L_MAX-interval)*sizeof(int));

	int **RI_array = (int **)malloc((L_MAX-interval)*sizeof(int*));
	int *dst_RI = (int *)malloc(NoC*(L_MAX-interval)*sizeof(int));

	int **OI_array = (int **)malloc((L_MAX-interval)*sizeof(int*));
	int *dst_OI = (int *)malloc((NoC)*(L_MAX-interval)*sizeof(int));

	int **RL_S_array = (int **)malloc((L_MAX-interval)*sizeof(int*));
	int *dst_RL_S = (int *)malloc(NoC*(L_MAX-interval)*sizeof(int));


	FLOAT **OFF_array = (FLOAT **)malloc((L_MAX-interval)*sizeof(FLOAT*));
	FLOAT *dst_OFF = (FLOAT *)malloc(NoC*(L_MAX-interval)*sizeof(FLOAT));

	FLOAT ***SCORE_array = (FLOAT ***)malloc((L_MAX-interval)*sizeof(FLOAT **));
	FLOAT **sub_dst_SCORE = (FLOAT **)malloc(NoC*(L_MAX-interval)*sizeof(FLOAT*));

	uintptr_t pointer_RL = (uintptr_t)dst_RL;
	uintptr_t pointer_RI = (uintptr_t)dst_RI;
	uintptr_t pointer_OI = (uintptr_t)dst_OI;
	uintptr_t pointer_RL_S = (uintptr_t)dst_RL_S;
	uintptr_t pointer_OFF = (uintptr_t)dst_OFF;
	uintptr_t pointer_SCORE = (uintptr_t)sub_dst_SCORE;
	for (int level=interval; level<L_MAX; level++) {

		int L=level-interval;

		RL_array[L] = (int *)pointer_RL;
		pointer_RL += (uintptr_t)NoC*sizeof(int);

		RI_array[L] = (int *)pointer_RI;
		pointer_RI += (uintptr_t)NoC*sizeof(int);

		OI_array[L] = (int *)pointer_OI;
		pointer_OI += (uintptr_t)NoC*sizeof(int);

		RL_S_array[L] = (int *)pointer_RL_S;
		pointer_RL_S += (uintptr_t)NoC*sizeof(int);

		OFF_array[L] = (FLOAT *)pointer_OFF;
		pointer_OFF += (uintptr_t)NoC*sizeof(FLOAT);

		SCORE_array[L] = (FLOAT **)pointer_SCORE;
		pointer_SCORE += (uintptr_t)NoC*sizeof(FLOAT*);
	}

	int sum_RL_S = 0;
	int sum_SNJ = 0;
	/* prepare for parallel execution */
	for(int level=interval; level<L_MAX; level++) {
		int L = level - interval;

		if(feature_size[level*2]+2*pady<MO->MI->max_Y ||(feature_size[level*2+1]+2*padx<MO->MI->max_X))
		{
			continue;
		}

		for(int j=0; j<NoC; j++) {

			/* root score + offset */
			RL_array[L][j] = rm_size_array[level][j*2]*rm_size_array[level][j*2+1];  //length of root-matching
			RI_array[L][j] = MO->MI->ridx[j];  //root-index
			OI_array[L][j] =  MO->MI->oidx[j];  //offset-index
			RL_S_array[L][j] =sizeof(FLOAT)*RL_array[L][j];


			OFF_array[L][j] = MO->MI->offw[RI_array[L][j]];  //offset information


			/* search max values */
			max_RL_S = (max_RL_S < RL_S_array[L][j]) ? RL_S_array[L][j] : max_RL_S;
			max_numpart = (max_numpart < numpart[j]) ? numpart[j] : max_numpart;
		}
	}

	sum_RL_S = max_RL_S*NoC*(L_MAX-interval);

	/* root matching size */
	sum_SNJ = sizeof(int*)*max_numpart*NoC*(L_MAX-interval);

	/* consolidated allocation for SCORE_array and distribute region */
	FLOAT *dst_SCORE = (FLOAT *)malloc(sum_RL_S);
	pointer_SCORE = (uintptr_t)dst_SCORE;
	for(int level=interval; level<L_MAX; level++) {
		int L = level - interval;

		if(feature_size[level*2]+2*pady<MO->MI->max_Y ||(feature_size[level*2+1]+2*padx<MO->MI->max_X))
		{
			continue;
		}

		for(int j=0; j<NoC; j++) {
			SCORE_array[L][j] = (FLOAT *)pointer_SCORE;
			pointer_SCORE += (uintptr_t)max_RL_S;
		}
	}

	/* add offset */
	for(int level=interval; level<L_MAX; level++) {
		int L = level - interval;

		if(feature_size[level*2]+2*pady<MO->MI->max_Y ||(feature_size[level*2+1]+2*padx<MO->MI->max_X))
		{
			continue;
		}

		for(int j=0; j<NoC; j++) {
			memcpy(SCORE_array[L][j], rootmatch[level][j], RL_S_array[L][j]);
			FLOAT *SC_S = SCORE_array[L][j];
			FLOAT *SC_E = SCORE_array[L][j]+RL_array[L][j];
			while(SC_S<SC_E) *(SC_S++)+=OFF_array[L][j];
		}
	}

	/* anchor matrix */  // consolidated allocation
	int ***ax_array = (int ***)malloc((L_MAX-interval)*sizeof(int **));
	int **sub_dst_ax = (int **)malloc(NoC*(L_MAX-interval)*sizeof(int *));
	int *dst_ax = (int *)malloc(sum_SNJ);

	int ***ay_array = (int ***)malloc((L_MAX-interval)*sizeof(int **));
	int **sub_dst_ay = (int **)malloc(NoC*(L_MAX-interval)*sizeof(int *));
	int *dst_ay = (int *)malloc(sum_SNJ);

	/* boudary index */  // consolidated allocation
	int ****Ix_array =(int ****)malloc((L_MAX-interval)*sizeof(int ***));
	int ***sub_dst_Ix = (int ***)malloc(NoC*(L_MAX-interval)*sizeof(int **));
	int **dst_Ix = (int **)malloc(sum_SNJ);

	int ****Iy_array = (int ****)malloc((L_MAX-interval)*sizeof(int ***));
	int ***sub_dst_Iy = (int ***)malloc(NoC*(L_MAX-interval)*sizeof(int **));
	int **dst_Iy = (int **)malloc(sum_SNJ);

	/* distribute region */
	uintptr_t pointer_ax = (uintptr_t)sub_dst_ax;
	uintptr_t pointer_ay = (uintptr_t)sub_dst_ay;
	uintptr_t pointer_Ix = (uintptr_t)sub_dst_Ix;
	uintptr_t pointer_Iy = (uintptr_t)sub_dst_Iy;
	for(int level=interval; level<L_MAX; level++) {
		int L = level - interval;

		if(feature_size[level*2]+2*pady<MO->MI->max_Y ||(feature_size[level*2+1]+2*padx<MO->MI->max_X))
		{
			continue;
		}

		ax_array[L] = (int **)pointer_ax;
		pointer_ax += (uintptr_t)(NoC*sizeof(int*));

		ay_array[L] = (int **)pointer_ay;
		pointer_ay += (uintptr_t)(NoC*sizeof(int*));

		Ix_array[L] = (int ***)pointer_Ix;
		pointer_Ix += (uintptr_t)(NoC*sizeof(int**));

		Iy_array[L] = (int ***)pointer_Iy;
		pointer_Iy += (uintptr_t)(NoC*sizeof(int**));
	}

	pointer_ax = (uintptr_t)dst_ax;
	pointer_ay = (uintptr_t)dst_ay;
	pointer_Ix = (uintptr_t)dst_Ix;
	pointer_Iy = (uintptr_t)dst_Iy;
	for(int level=interval; level<L_MAX; level++) {
		int L = level - interval;

		if(feature_size[level*2]+2*pady<MO->MI->max_Y ||(feature_size[level*2+1]+2*padx<MO->MI->max_X))
		{
			continue;
		}

		for(int j=0; j<NoC; j++) {
			uintptr_t pointer_offset = sizeof(int*)*max_numpart;

			ax_array[L][j] = (int *)pointer_ax;
			pointer_ax += pointer_offset;

			ay_array[L][j] = (int *)pointer_ay;
			pointer_ay += pointer_offset;

			Ix_array[L][j] = (int **)pointer_Ix;
			pointer_Ix += pointer_offset;

			Iy_array[L][j] = (int **)pointer_Iy;
			pointer_Iy += pointer_offset;
		}
	}

	/* add parts */
	if(NoP>0)
        {
		/* arrays to store temporary loop variables */
		int tmp_array_size = 0;
		for(int level=interval; level<L_MAX; level++) {
			if(feature_size[level*2]+2*pady<MO->MI->max_Y ||(feature_size[level*2+1]+2*padx<MO->MI->max_X))
			{
				continue;
			}

			for(int j=0; j<NoC; j++) {
				tmp_array_size += max_numpart*sizeof(int);
			}
		}

		int ***DIDX_array = (int ***)malloc((L_MAX-interval)*sizeof(int**));
		int **sub_dst_DIDX = (int **)malloc(NoC*(L_MAX-interval)*sizeof(int*));
		int *dst_DIDX = (int *)malloc(tmp_array_size);


		int ***DID_4_array = (int ***)malloc((L_MAX-interval)*sizeof(int **));
		int **sub_dst_DID_4 = (int **)malloc(NoC*(L_MAX-interval)*sizeof(int*));
		int *dst_DID_4;
		res = cuMemHostAlloc((void **)&dst_DID_4, tmp_array_size, CU_MEMHOSTALLOC_DEVICEMAP);
		if(res != CUDA_SUCCESS) {
			printf("cuMemHostAlloc(dst_DID_4) failed: res = %s\n", cuda_response_to_string(res));
			exit(1);
		}


		int ***PIDX_array = (int ***)malloc((L_MAX-interval)*sizeof(int **));
		int **sub_dst_PIDX = (int **)malloc(NoC*(L_MAX-interval)*sizeof(int*));
		int *dst_PIDX;
		res = cuMemHostAlloc((void **)&dst_PIDX, tmp_array_size, CU_MEMHOSTALLOC_DEVICEMAP);
		if(res != CUDA_SUCCESS) {
			printf("cuMemHostAlloc(dst_PIDX) failed: res = %s\n", cuda_response_to_string(res));
			exit(1);
		}

		/* distribute consolidated region */
		uintptr_t pointer_DIDX = (uintptr_t)sub_dst_DIDX;
		uintptr_t pointer_DID_4 = (uintptr_t)sub_dst_DID_4;
		uintptr_t pointer_PIDX = (uintptr_t)sub_dst_PIDX;
		for(int level=interval; level<L_MAX; level++) {
			int L = level - interval;

			if(feature_size[level*2]+2*pady<MO->MI->max_Y ||(feature_size[level*2+1]+2*padx<MO->MI->max_X))
			{
				continue;
			}

			DIDX_array[L] = (int **)pointer_DIDX;
			pointer_DIDX += (uintptr_t)(NoC*sizeof(int*));

			DID_4_array[L] = (int **)pointer_DID_4;
			pointer_DID_4 += (uintptr_t)(NoC*sizeof(int*));

			PIDX_array[L] = (int **)pointer_PIDX;
			pointer_PIDX += (uintptr_t)(NoC*sizeof(int*));
		}

		pointer_DIDX = (uintptr_t)dst_DIDX;
		pointer_DID_4 = (uintptr_t)dst_DID_4;
		pointer_PIDX = (uintptr_t)dst_PIDX;
		for(int level=interval; level<L_MAX; level++) {
			int L = level - interval;

			if(feature_size[level*2]+2*pady<MO->MI->max_Y ||(feature_size[level*2+1]+2*padx<MO->MI->max_X)) {
				continue;
			}

			for(int j=0; j<NoC; j++) {
				uintptr_t pointer_offset = (uintptr_t)(max_numpart*sizeof(int));

				DIDX_array[L][j] = (int *)pointer_DIDX;
				pointer_DIDX += pointer_offset;

				DID_4_array[L][j] = (int *)pointer_DID_4;
				pointer_DID_4 += pointer_offset;

				PIDX_array[L][j] = (int *)pointer_PIDX;
				pointer_PIDX += pointer_offset;
			}
		}

		/* prepare for parallel execution */
		int sum_size_index_matrix = 0;
		for(int level=interval; level<L_MAX; level++) {
			int L = level - interval;

			if(feature_size[level*2]+2*pady<MO->MI->max_Y ||(feature_size[level*2+1]+2*padx<MO->MI->max_X)) {
				continue;
			}

			for(int j=0; j<NoC; j++) {
				for (int k=0;k<numpart[j];k++) {
					/* assign values to each element */
					DIDX_array[L][j][k] = MO->MI->didx[j][k];
					DID_4_array[L][j][k] = DIDX_array[L][j][k]*4;
					PIDX_array[L][j][k] = MO->MI->pidx[j][k];

					/* anchor */
					ax_array[L][j][k] = MO->MI->anchor[DIDX_array[L][j][k]*2]+1;
					ay_array[L][j][k] = MO->MI->anchor[DIDX_array[L][j][k]*2+1]+1;

					int PSSIZE[2] ={pm_size_array[L][PIDX_array[L][j][k]*2], pm_size_array[L][PIDX_array[L][j][k]*2+1]}; // size of C

					/* index matrix */
					sum_size_index_matrix += sizeof(int)*PSSIZE[0]*PSSIZE[1];
				}
			}
		}

		int *dst_Ix_kk = (int *)malloc(sum_size_index_matrix);
		int *dst_Iy_kk = (int *)malloc(sum_size_index_matrix);
		uintptr_t pointer_Ix_kk = (uintptr_t)dst_Ix_kk;
		uintptr_t pointer_Iy_kk = (uintptr_t)dst_Iy_kk;
		for(int level=interval; level<L_MAX; level++) {
			int L = level - interval;

			if(feature_size[level*2]+2*pady<MO->MI->max_Y ||(feature_size[level*2+1]+2*padx<MO->MI->max_X))
			{
				continue;
			}

			for(int j=0; j<NoC; j++) {
				for (int k=0;k<numpart[j];k++) {
					int PSSIZE[2] ={pm_size_array[L][PIDX_array[L][j][k]*2], pm_size_array[L][PIDX_array[L][j][k]*2+1]}; // size of C

					Ix_array[L][j][k] = (int *)pointer_Ix_kk;
					Iy_array[L][j][k] = (int *)pointer_Iy_kk;

					pointer_Ix_kk += (uintptr_t)(sizeof(int)*PSSIZE[0]*PSSIZE[1]);
					pointer_Iy_kk += (uintptr_t)(sizeof(int)*PSSIZE[0]*PSSIZE[1]);
				}
			}
		}

		gettimeofday(&tv_dt_start, nullptr);
		FLOAT ****M_array = dt_GPU(
			Ix_array,      // int ****Ix_array
			Iy_array,      // int ****Iy_array
			PIDX_array,    // int ***PIDX_array
			pm_size_array, // int **size_array
			NoP,           // int NoP
			numpart,       // int *numpart
			NoC,           // int NoC
			interval,      // int interval
			L_MAX,         // int L_MAX
			feature_size,         // int *feature_size,
			padx,          // int padx,
			pady,          // int pady,
			MO->MI->max_X, // int max_X
			MO->MI->max_Y, // int max_Y
			MO->MI->def, // FLOAT *def
			tmp_array_size, // int tmp_array_size
			dst_PIDX, // int *dst_PIDX
			dst_DID_4 // int *DID_4
			);
		gettimeofday(&tv_dt_end, nullptr);
		tvsub(&tv_dt_end, &tv_dt_start, &tv);
		time_dt += tv.tv_sec * 1000.0 + (float)tv.tv_usec / 1000.0;

		/* add part score */
		for(int level=interval; level<L_MAX; level++){
			int L = level - interval;

			if(feature_size[level*2]+2*pady<MO->MI->max_Y ||(feature_size[level*2+1]+2*padx<MO->MI->max_X))
			{
				continue;
			}

			for(int j=0; j<NoC; j++) {
				for(int k=0; k<numpart[j]; k++) {
					int PSSIZE[2] ={pm_size_array[L][PIDX_array[L][j][k]*2],
							pm_size_array[L][PIDX_array[L][j][k]*2+1]}; // Size of C
					int R_S[2]={rm_size_array[level][j*2], rm_size_array[level][j*2+1]};

					dpm_ttic_add_part_calculation(SCORE_array[L][j], M_array[L][j][k], R_S,
								      PSSIZE, ax_array[L][j][k], ay_array[L][j][k]);
				}
			}
		}

		s_free(M_array[0][0][0]);
		s_free(M_array[0][0]);
		s_free(M_array[0]);
		s_free(M_array);

		/* free temporary arrays */
		free(dst_DIDX);
		free(sub_dst_DIDX);
		free(DIDX_array);

		res = cuMemFreeHost(dst_DID_4);
		if(res != CUDA_SUCCESS) {
			printf("cuMemFreeHost(dst_DID_4) failed: res = %s\n", cuda_response_to_string(res));
			exit(1);
		}
		free(sub_dst_DID_4);
		free(DID_4_array);

		res = cuMemFreeHost(dst_PIDX);
		if(res != CUDA_SUCCESS) {
			printf("cuMemFreeHost(dst_PIDX) failed: res = %s\n", cuda_response_to_string(res));
			exit(1);
		}

		free(sub_dst_PIDX);
		free(PIDX_array);

		res = cuCtxSetCurrent(ctx[0]);
		if(res != CUDA_SUCCESS) {
			printf("cuCtxSetCurrent(ctx[0]) failed: res = %s\n", cuda_response_to_string(res));
			exit(1);
		}
        } // start from if(NoP>0)

	/* combine root and part score and detect boundary box for each-component */

	FLOAT *scale_array = (FLOAT *)malloc((L_MAX-interval)*sizeof(FLOAT));
	for(int level=interval; level<L_MAX; level++) {
		int L = level - interval;

		if(feature_size[level*2]+2*pady<MO->MI->max_Y ||(feature_size[level*2+1]+2*padx<MO->MI->max_X))
		{
			Tboxes[count]=nullptr;
			count++;
			continue;
		}

		scale_array[L] = (FLOAT)sbin/scales[level];
	}

	for (int level=interval; level<L_MAX; level++)  // feature's loop(A's loop) 1level 1picture
        {
		/* parameters (related for level) */
		int L=level-interval;
		/* matched score size matrix */
		FLOAT scale=(FLOAT)sbin/scales[level];

		/* loop conditon */
		if(feature_size[level*2]+2*pady<MO->MI->max_Y ||(feature_size[level*2+1]+2*padx<MO->MI->max_X)) {
			Tboxes[count]=nullptr;
			count++;
			continue;
		}

		/* calculate accumulated score */
		gettimeofday(&tv_calc_a_score_start, nullptr);

		calc_a_score_GPU(
			acc_score,              // FLOAT *ac_score
			SCORE_array[L],       // FLOAT **score
			rm_size_array[level], // int *ssize_start
			MO->MI,               // Model_info *MI
			scale,                // FLOAT scale
			RL_S_array[L],        // int *size_score_array
			NoC                   // int NoC
			);

		gettimeofday(&tv_calc_a_score_end, nullptr);
		tvsub(&tv_calc_a_score_end, &tv_calc_a_score_start, &tv);
		time_calc_a_score += tv.tv_sec * 1000.0 + (float)tv.tv_usec / 1000.0;

		for(int j = 0; j <NoC; j++) {
			int R_S[2]={rm_size_array[level][j*2], rm_size_array[level][j*2+1]};

			/* get all good matches */
			int GMN;
			int *GMPC = get_gmpc(SCORE_array[L][j],thresh,R_S,&GMN);
			int RSIZE[2]={MO->MI->rsize[j*2], MO->MI->rsize[j*2+1]};

			int GL = (numpart[j]+1)*4+3;  //31

			/* detected box coordinate(current level) */
			FLOAT *t_boxes = (FLOAT*)calloc(GMN*GL,sizeof(FLOAT));

			gettimeofday(&tv_box_start, nullptr);

			// NO NEED TO USE GPU 
			for(int k = 0;k < GMN;k++) {
				FLOAT *P_temp = t_boxes+GL*k;
				int y = GMPC[2*k];
				int x = GMPC[2*k+1];

				/* calculate root box coordinate */
				FLOAT *RB =rootbox(x,y,scale,padx,pady,RSIZE);
				memcpy(P_temp, RB,sizeof(FLOAT)*4);
				s_free(RB);
				P_temp+=4;

				for(int pp=0;pp<numpart[j];pp++) {
					int PBSIZE[2]={psize[j][pp*2], psize[j][pp*2+1]};
					int Isize[2]={pm_size_array[L][MO->MI->pidx[j][pp]*2], pm_size_array[L][MO->MI->pidx[j][pp]*2+1]};

					/* calculate part box coordinate */
					FLOAT *PB = partbox(x,y,ax_array[L][j][pp],ay_array[L][j][pp],scale,padx,pady,PBSIZE,Ix_array[L][j][pp],Iy_array[L][j][pp],Isize);
					memcpy(P_temp, PB,sizeof(FLOAT)*4);
					P_temp+=4;
					s_free(PB);
				}
				/* component number and score */
				*(P_temp++)=(FLOAT)j; //component number
				*(P_temp++)=SCORE_array[L][j][x*R_S[0]+y]; //score of good match
				*P_temp = scale;
			}

			//  NO NEED TO USE GPU
			gettimeofday(&tv_box_end, nullptr);
			tvsub(&tv_box_end, &tv_box_start, &tv);
			time_box += tv.tv_sec * 1000.0 + (float)tv.tv_usec / 1000.0;

			/* save box information */
			if (GMN > 0)
				Tboxes[count] = t_boxes;
			else
				Tboxes[count] = nullptr;

			b_nums[count]=GMN;
			count++;
			detected_boxes+=GMN;			//number of detected box

			/* release */
			s_free(GMPC);
		}
		////numcom
        }
	////level

	/* free temporary arrays */
	free(dst_RL);
	free(RL_array);

	free(dst_RI);
	free(RI_array);

	free(dst_OI);
	free(OI_array);

	free(dst_RL_S);
	free(RL_S_array);

	free(dst_OFF);
	free(OFF_array);

	free(dst_SCORE);
	free(sub_dst_SCORE);
	free(SCORE_array);

	free(dst_ax);
	free(sub_dst_ax);
	free(ax_array);

	free(dst_ay);
	free(sub_dst_ay);
	free(ay_array);

	free(Ix_array[0][0][0]);
	free(dst_Ix);
	free(sub_dst_Ix);
	free(Ix_array);

	free(Iy_array[0][0][0]);
	free(dst_Iy);
	free(sub_dst_Iy);
	free(Iy_array);

	free(scale_array);

	gettimeofday(&tv_nucom_end, nullptr);

#ifdef PRINT_INFO
	printf("root SCORE : %f\n", time_root_score);
	printf("part SCORE : %f\n", time_part_score);
	printf("dt  : %f\n", time_dt);
	printf("calc_a_score : %f\n", time_calc_a_score);
#endif
	res = cuCtxSetCurrent(ctx[0]);
	if(res != CUDA_SUCCESS) {
		printf("cuCtxSetCurrent(ctx[0]) failed: res = %s\n",cuda_response_to_string(res));
		exit(1);
	}

	/* free memory regions */
	res = cuMemFreeHost((void *)featp2[0]);
	if(res != CUDA_SUCCESS) {
		printf("cuMemFreeHost(featp2[0]) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}

	s_free(featp2);

	res = cuMemFreeHost((void *)rootmatch[interval][0]);
	if(res != CUDA_SUCCESS) {
		printf("cuMemFreeHost(rootmatch[0][0]) failed: res = %s\n", cuda_response_to_string(res));
		exit(1);
	}
	s_free(rootmatch[0]);
	s_free(rootmatch);

	if (partmatch != nullptr) {
		res = cuMemFreeHost((void *)partmatch[0][0]);
		if(res != CUDA_SUCCESS) {
			printf("cuMemFreeHost(partmatch[0][0]) failed: res = %s\n", cuda_response_to_string(res));
			exit(1);
		}

		s_free(partmatch[0]);
		s_free(partmatch);

		s_free(new_PADsize);
	}

	/* release */
	s_free(rm_size_array[0]);
	s_free(rm_size_array);
	s_free(pm_size_array[0]);
	s_free(pm_size_array);

	/* Output boundary-box coorinate information */
	int GL=(numpart[0]+1)*4+3;
	FLOAT *boxes=(FLOAT*)calloc(detected_boxes*GL,sizeof(FLOAT));		//box coordinate information(Temp)

	FLOAT *T1 = boxes;
	for(int i = 0; i < LofFeat; i++) {
		int num_t = b_nums[i]*GL;
		if(num_t > 0) {
			FLOAT *T2 = Tboxes[i];
			//memcpy_s(T1,sizeof(FLOAT)*num_t,T2,sizeof(FLOAT)*num_t);
			memcpy(T1, T2,sizeof(FLOAT)*num_t);
			T1 += num_t;
		}
	}

	FLOAT abs_threshold = abs(thresh);

	/* accumulated score calculation */
	FLOAT max_score = 0.0;

	/* add offset to accumulated score */
	for(int i = 0; i < MO->MI->IM_HEIGHT*MO->MI->IM_WIDTH; i++) {
		if (acc_score[i] < thresh) {
			acc_score[i] = 0.0;
		} else {
			acc_score[i] += abs_threshold;

			if (acc_score[i] > max_score)
				max_score = acc_score[i];
		}
	}

	/* normalization */
	if (max_score > 0.0) {
		FLOAT ac_ratio = 1.0 / max_score;

		for (int i = 0; i < MO->MI->IM_HEIGHT*MO->MI->IM_WIDTH; i++) {
			acc_score[i] *= ac_ratio;
		}
	}

	/* release */
	free_boxes(Tboxes,LofFeat);
	s_free(b_nums);

	/* output result */
	*detected_count = detected_boxes;
	return boxes;
}
Ejemplo n.º 9
0
cl_int
pocl_cuda_alloc_mem_obj (cl_device_id device, cl_mem mem_obj, void *host_ptr)
{
  cuCtxSetCurrent (((pocl_cuda_device_data_t *)device->data)->context);

  CUresult result;
  void *b = NULL;

  /* if memory for this global memory is not yet allocated -> do it */
  if (mem_obj->device_ptrs[device->global_mem_id].mem_ptr == NULL)
    {
      cl_mem_flags flags = mem_obj->flags;

      if (flags & CL_MEM_USE_HOST_PTR)
        {
#if defined __arm__
          // cuMemHostRegister is not supported on ARN
          // Allocate device memory and perform explicit copies
          // before and after running a kernel
          result = cuMemAlloc ((CUdeviceptr *)&b, mem_obj->size);
          CUDA_CHECK (result, "cuMemAlloc");
#else
          result = cuMemHostRegister (host_ptr, mem_obj->size,
                                      CU_MEMHOSTREGISTER_DEVICEMAP);
          if (result != CUDA_SUCCESS
              && result != CUDA_ERROR_HOST_MEMORY_ALREADY_REGISTERED)
            CUDA_CHECK (result, "cuMemHostRegister");
          result = cuMemHostGetDevicePointer ((CUdeviceptr *)&b, host_ptr, 0);
          CUDA_CHECK (result, "cuMemHostGetDevicePointer");
#endif
        }
      else if (flags & CL_MEM_ALLOC_HOST_PTR)
        {
          result = cuMemHostAlloc (&mem_obj->mem_host_ptr, mem_obj->size,
                                   CU_MEMHOSTREGISTER_DEVICEMAP);
          CUDA_CHECK (result, "cuMemHostAlloc");
          result = cuMemHostGetDevicePointer ((CUdeviceptr *)&b,
                                              mem_obj->mem_host_ptr, 0);
          CUDA_CHECK (result, "cuMemHostGetDevicePointer");
        }
      else
        {
          result = cuMemAlloc ((CUdeviceptr *)&b, mem_obj->size);
          if (result != CUDA_SUCCESS)
            {
              const char *err;
              cuGetErrorName (result, &err);
              POCL_MSG_PRINT2 (__FUNCTION__, __LINE__,
                               "-> Failed to allocate memory: %s\n", err);
              return CL_MEM_OBJECT_ALLOCATION_FAILURE;
            }
        }

      if (flags & CL_MEM_COPY_HOST_PTR)
        {
          result = cuMemcpyHtoD ((CUdeviceptr)b, host_ptr, mem_obj->size);
          CUDA_CHECK (result, "cuMemcpyHtoD");
        }

      mem_obj->device_ptrs[device->global_mem_id].mem_ptr = b;
      mem_obj->device_ptrs[device->global_mem_id].global_mem_id
          = device->global_mem_id;
    }

  /* copy already allocated global mem info to devices own slot */
  mem_obj->device_ptrs[device->dev_id]
      = mem_obj->device_ptrs[device->global_mem_id];

  return CL_SUCCESS;
}
Ejemplo n.º 10
0
/*
 * Class:     edu_syr_pcpratts_rootbeer_runtime2_cuda_CudaRuntime2
 * Method:    findReserveMem
 * Signature: ()I
 */
JNIEXPORT jlong JNICALL Java_edu_syr_pcpratts_rootbeer_runtime2_cuda_CudaRuntime2_findReserveMem
  (JNIEnv * env, jobject this_ref, jint max_blocks_per_proc, jint max_threads_per_block)
{
  size_t to_space_size;
  size_t temp_size;
  int status;
  int deviceCount = 0;
  jlong prev_i;
  jlong i;
  size_t f_mem;
  size_t t_mem;
  jint num_blocks;

  status = cuInit(0);
  CHECK_STATUS(env,"error in cuInit",status)

  printf("automatically determining CUDA reserve space...\n");
  
  to_space_size = initContext(env, max_blocks_per_proc, max_threads_per_block);

  //space for 100 types in the scene
  classMemSize = sizeof(jint)*100;

  num_blocks = numMultiProcessors * max_threads_per_block * max_blocks_per_proc;
  
  gc_space_size = 1024;
  to_space_size -= (num_blocks * sizeof(jlong));
  to_space_size -= (num_blocks * sizeof(jlong));
  to_space_size -= gc_space_size;
  to_space_size -= classMemSize;
  
  for(i = 1024L*1024L; i < to_space_size; i += 100L*1024L*1024L){
    temp_size = to_space_size - i;
  
    printf("attempting allocation with temp_size: %lu to_space_size: %lu i: %ld\n", temp_size, to_space_size, i);
 
    status = cuMemHostAlloc(&toSpace, temp_size, 0);  
    if(status != CUDA_SUCCESS){
	    cuCtxDestroy(cuContext);
	    initContext(env, max_blocks_per_proc, max_threads_per_block);
      continue;
    }
    
    status = cuMemAlloc(&gpuToSpace, temp_size);
    if(status != CUDA_SUCCESS){
	    cuCtxDestroy(cuContext);
	    initContext(env, max_blocks_per_proc, max_threads_per_block);
      continue;
    } 

    status = cuMemAlloc(&gpuClassMemory, classMemSize);
    if(status != CUDA_SUCCESS){
	    cuCtxDestroy(cuContext);
	    initContext(env, max_blocks_per_proc, max_threads_per_block);
      continue;
    } 

    status = cuMemHostAlloc(&handlesMemory, num_blocks * sizeof(jlong), CU_MEMHOSTALLOC_WRITECOMBINED); 
    if(status != CUDA_SUCCESS){
	    cuCtxDestroy(cuContext);
	    initContext(env, max_blocks_per_proc, max_threads_per_block);
      continue;
    } 

    status = cuMemAlloc(&gpuHandlesMemory, num_blocks * sizeof(jlong)); 
    if(status != CUDA_SUCCESS){
	    cuCtxDestroy(cuContext);
	    initContext(env, max_blocks_per_proc, max_threads_per_block);
      continue;
    } 

    status = cuMemHostAlloc(&exceptionsMemory, num_blocks * sizeof(jlong), 0); 
    if(status != CUDA_SUCCESS){
	    cuCtxDestroy(cuContext);
	    initContext(env, max_blocks_per_proc, max_threads_per_block);
      continue;
    } 

    status = cuMemAlloc(&gpuExceptionsMemory, num_blocks * sizeof(jlong)); 
    if(status != CUDA_SUCCESS){
	    cuCtxDestroy(cuContext);
	    initContext(env, max_blocks_per_proc, max_threads_per_block);
      continue;
    } 

    status = cuMemAlloc(&gcInfoSpace, gc_space_size);  
    if(status != CUDA_SUCCESS){
	    cuCtxDestroy(cuContext);
	    initContext(env, max_blocks_per_proc, max_threads_per_block);
      continue;
    } 

    status = cuMemAlloc(&gpuHeapEndPtr, 8);
    if(status != CUDA_SUCCESS){
	    cuCtxDestroy(cuContext);
	    initContext(env, max_blocks_per_proc, max_threads_per_block);
      continue;
    } 

    status = cuMemAlloc(&gpuBufferSize, 8);
    if(status != CUDA_SUCCESS){
	    cuCtxDestroy(cuContext);
	    initContext(env, max_blocks_per_proc, max_threads_per_block);
      continue;
    } 

    //done, free everything
    cuMemFree(gpuToSpace);
    cuMemFree(gpuClassMemory);
    cuMemFree(gpuHandlesMemory);
    cuMemFree(gpuExceptionsMemory);
    cuMemFree(gcInfoSpace);
    cuMemFree(gpuHeapEndPtr);
    cuMemFree(gpuBufferSize);

	  cuMemFreeHost(toSpace);
	  cuMemFreeHost(handlesMemory);
	  cuMemFreeHost(exceptionsMemory);

    return i;
  }
  throw_cuda_errror_exception(env, "unable to find enough space using CUDA", 0); 
  return 0;
}
Ejemplo n.º 11
0
void initDevice(JNIEnv * env, jobject this_ref, jint max_blocks_per_proc, jint max_threads_per_block, jlong free_space)
{          
  int status;
  jint num_blocks;
  int deviceCount = 0;
  size_t f_mem;
  size_t t_mem;
  size_t to_space_size;
  textureMemSize = 1;

  status = cuDeviceGetCount(&deviceCount);
  CHECK_STATUS(env,"error in cuDeviceGetCount",status)

  getBestDevice(env);
  
  status = cuCtxCreate(&cuContext, CU_CTX_MAP_HOST, cuDevice);  
  CHECK_STATUS(env,"error in cuCtxCreate",status)
  
  status = cuMemGetInfo (&f_mem, &t_mem);
  CHECK_STATUS(env,"error in cuMemGetInfo",status)
          
  to_space_size = f_mem;
  
  num_blocks = numMultiProcessors * max_threads_per_block * max_blocks_per_proc;
  
#if DEBUG

  printf("Memory: %i(MB)/%i(MB) (Free/Total)\n",f_mem/1024/1024, t_mem/1024/1024);
  
  printf("num_blocks = %i\n",num_blocks);
  printf("numMultiProcessors = %i\n",numMultiProcessors);
  printf("max_threads_per_block = %i\n",max_threads_per_block);
  printf("max_blocks_per_proc = %i\n",max_blocks_per_proc);
  fflush(stdout);
#endif

  //space for 100 types in the scene
  classMemSize = sizeof(jint)*100;
  
  gc_space_size = 1024;
  to_space_size -= (num_blocks * sizeof(jlong));
  to_space_size -= (num_blocks * sizeof(jlong));
  to_space_size -= gc_space_size;
  to_space_size -= free_space;
  to_space_size -= classMemSize;

  //to_space_size -= textureMemSize;
  bufferSize = to_space_size;

  status = cuMemHostAlloc(&toSpace, to_space_size, 0);  
  CHECK_STATUS(env,"toSpace memory allocation failed",status)
    
  status = cuMemAlloc(&gpuToSpace, to_space_size);
  CHECK_STATUS(env,"gpuToSpace memory allocation failed",status)
    
  status = cuMemAlloc(&gpuClassMemory, classMemSize);
  CHECK_STATUS(env,"gpuClassMemory memory allocation failed",status)
  
/*
  status = cuMemHostAlloc(&textureMemory, textureMemSize, 0);  
  if (CUDA_SUCCESS != status) 
  {
    printf("error in cuMemHostAlloc textureMemory %d\n", status);
  }

  status = cuMemAlloc(&gpuTexture, textureMemSize);
  if (CUDA_SUCCESS != status) 
  {
    printf("error in cuMemAlloc gpuTexture %d\n", status);
  }
*/

  status = cuMemHostAlloc(&handlesMemory, num_blocks * sizeof(jlong), CU_MEMHOSTALLOC_WRITECOMBINED); 
  CHECK_STATUS(env,"handlesMemory memory allocation failed",status)

  status = cuMemAlloc(&gpuHandlesMemory, num_blocks * sizeof(jlong)); 
  CHECK_STATUS(env,"gpuHandlesMemory memory allocation failed",status)

  status = cuMemHostAlloc(&exceptionsMemory, num_blocks * sizeof(jlong), 0); 
  CHECK_STATUS(env,"exceptionsMemory memory allocation failed",status)

  status = cuMemAlloc(&gpuExceptionsMemory, num_blocks * sizeof(jlong)); 
  CHECK_STATUS(env,"gpuExceptionsMemory memory allocation failed",status)

  status = cuMemAlloc(&gcInfoSpace, gc_space_size);  
  CHECK_STATUS(env,"gcInfoSpace memory allocation failed",status)

  status = cuMemAlloc(&gpuHeapEndPtr, 8);
  CHECK_STATUS(env,"gpuHeapEndPtr memory allocation failed",status)

  status = cuMemAlloc(&gpuBufferSize, 8);
  CHECK_STATUS(env,"gpuBufferSize memory allocation failed",status)

  thisRefClass = (*env)->GetObjectClass(env, this_ref);
  setLongField(env, this_ref, "m_ToSpaceAddr", (jlong) toSpace);
  setLongField(env, this_ref, "m_GpuToSpaceAddr", (jlong) gpuToSpace);
  setLongField(env, this_ref, "m_TextureAddr", (jlong) textureMemory);
  setLongField(env, this_ref, "m_GpuTextureAddr", (jlong) gpuTexture);
  setLongField(env, this_ref, "m_HandlesAddr", (jlong) handlesMemory);
  setLongField(env, this_ref, "m_GpuHandlesAddr", (jlong) gpuHandlesMemory);
  setLongField(env, this_ref, "m_ExceptionsHandlesAddr", (jlong) exceptionsMemory);
  setLongField(env, this_ref, "m_GpuExceptionsHandlesAddr", (jlong) gpuExceptionsMemory);
  setLongField(env, this_ref, "m_ToSpaceSize", (jlong) bufferSize);
  setLongField(env, this_ref, "m_MaxGridDim", (jlong) maxGridDim);
  setLongField(env, this_ref, "m_NumMultiProcessors", (jlong) numMultiProcessors);
}
Ejemplo n.º 12
0
/*
 * Class:     edu_syr_pcpratts_rootbeer_runtime2_cuda_CudaRuntime2
 * Method:    setup
 * Signature: ()V
 */
JNIEXPORT void JNICALL Java_edu_syr_pcpratts_rootbeer_runtime2_cuda_CudaRuntime2_setup
  (JNIEnv *env, jobject this_ref, jint max_blocks_per_proc, jint max_threads_per_block, jint free_space)
{
  int status;
  jint num_blocks;
  int deviceCount = 0;
  size_t f_mem;
  size_t t_mem;
  size_t to_space_size;
  //size_t free_space = 1530L*1024L*1024L;
  textureMemSize = 1;
  
  status = cuInit(0);
  if (CUDA_SUCCESS != status) 
  {
    printf("error in cuInit\n");
  }
  
  status = cuDeviceGetCount(&deviceCount);
  if (CUDA_SUCCESS != status) 
  {
    printf("error in cuDeviceGet\n");
  }

  getBestDevice();
  
  status = cuCtxCreate(&cuContext, CU_CTX_MAP_HOST, cuDevice);  
  if (CUDA_SUCCESS != status) 
  {
    printf("error in cuCtxCreate %d\n", status);
  }
  
  // ddb - not using this as this returns the total memory not the free memory
  //to_space_size = memSize();
  
  cuMemGetInfo(&f_mem, &t_mem);
  
  to_space_size = f_mem;
  
  num_blocks = numMultiProcessors * max_threads_per_block * max_blocks_per_proc;
  
#if DEBUG

  printf("Memory: %i(MB)/%i(MB) (Free/Total)\n",f_mem/1024/1024, t_mem/1024/1024);
  
  printf("num_blocks = %i\n",num_blocks);
  printf("numMultiProcessors = %i\n",numMultiProcessors);
  printf("max_threads_per_block = %i\n",max_threads_per_block);
  printf("max_blocks_per_proc = %i\n",max_blocks_per_proc);
  fflush(stdout);
#endif
  
  gc_space_size = 1024;
  to_space_size -= (num_blocks * sizeof(jlong));
  to_space_size -= (num_blocks * sizeof(jlong));
  to_space_size -= gc_space_size;
  to_space_size -= free_space;
  //to_space_size -= textureMemSize;
  bufferSize = to_space_size;

  status = cuMemHostAlloc(&toSpace, to_space_size, 0);  
  
  if (CUDA_SUCCESS != status) {
    throw_cuda_errror_exception(env, "toSpace memory allocation failed", status);
    return;
  }
  
  status = cuMemAlloc(&gpuToSpace, to_space_size);
  
  if (CUDA_SUCCESS != status) {
    throw_cuda_errror_exception(env, "gpuToSpace memory allocation failed", status);
    return;
  }
  
/*
  status = cuMemHostAlloc(&textureMemory, textureMemSize, 0);  
  if (CUDA_SUCCESS != status) 
  {
    printf("error in cuMemHostAlloc textureMemory %d\n", status);
  }

  status = cuMemAlloc(&gpuTexture, textureMemSize);
  if (CUDA_SUCCESS != status) 
  {
    printf("error in cuMemAlloc gpuTexture %d\n", status);
  }
*/
  status = cuMemHostAlloc(&handlesMemory, num_blocks * sizeof(jlong), CU_MEMHOSTALLOC_WRITECOMBINED); 
  
  if (CUDA_SUCCESS != status) {
    throw_cuda_errror_exception(env, "handlesMemory memory allocation failed", status);
    return;
  }

  status = cuMemAlloc(&gpuHandlesMemory, num_blocks * sizeof(jlong)); 
  
  if (CUDA_SUCCESS != status) {
    throw_cuda_errror_exception(env, "gpuHandlesMemory memory allocation failed", status);
    return;
  }

  status = cuMemHostAlloc(&exceptionsMemory, num_blocks * sizeof(jlong), 0); 
  
  if (CUDA_SUCCESS != status) {
    throw_cuda_errror_exception(env, "exceptionsMemory memory allocation failed", status);
    return;
  }

  status = cuMemAlloc(&gpuExceptionsMemory, num_blocks * sizeof(jlong)); 
 
  if (CUDA_SUCCESS != status) {
    throw_cuda_errror_exception(env, "gpuExceptionsMemory memory allocation failed", status);
    return;
  }

  status = cuMemAlloc(&gcInfoSpace, gc_space_size);  
  
  if (CUDA_SUCCESS != status) {
    throw_cuda_errror_exception(env, "gcInfoSpace memory allocation failed", status);
    return;
  }

  status = cuMemAlloc(&gpuHeapEndPtr, 8);
  
  if (CUDA_SUCCESS != status) {
    throw_cuda_errror_exception(env, "gpuHeapEndPtr memory allocation failed", status);
    return;
  }

  status = cuMemAlloc(&gpuBufferSize, 8);
  
  if (CUDA_SUCCESS != status) {
    throw_cuda_errror_exception(env, "gpuBufferSize memory allocation failed", status);
    return;
  }

  thisRefClass = (*env)->GetObjectClass(env, this_ref);
  setLongField(env, this_ref, "m_ToSpaceAddr", (jlong) toSpace);
  setLongField(env, this_ref, "m_GpuToSpaceAddr", (jlong) gpuToSpace);
  setLongField(env, this_ref, "m_TextureAddr", (jlong) textureMemory);
  setLongField(env, this_ref, "m_GpuTextureAddr", (jlong) gpuTexture);
  setLongField(env, this_ref, "m_HandlesAddr", (jlong) handlesMemory);
  setLongField(env, this_ref, "m_GpuHandlesAddr", (jlong) gpuHandlesMemory);
  setLongField(env, this_ref, "m_ExceptionsHandlesAddr", (jlong) exceptionsMemory);
  setLongField(env, this_ref, "m_GpuExceptionsHandlesAddr", (jlong) gpuExceptionsMemory);
  setLongField(env, this_ref, "m_ToSpaceSize", (jlong) bufferSize);
  setLongField(env, this_ref, "m_MaxGridDim", (jlong) maxGridDim);
  setLongField(env, this_ref, "m_NumMultiProcessors", (jlong) numMultiProcessors);
}
Ejemplo n.º 13
0
//load model basic information
Model_info * load_modelinfo(char *filename)
{	
  CUresult res;

  FILE *file;		//File
  errno_t err;	//err ( for fopen)
  Model_info *MI=(Model_info*)malloc(sizeof(Model_info));		//Model information
  
  //fopen	
  //if( (err = fopen_s( &file,filename, "r"))!=0 ) 
  if( (file=fopen(filename, "r"))==NULL ) 
    {
      printf("Model information file not found \n");
      exit(-1);
    }
  FLOAT t1,t2,t3,t4;	
  int b;

  //load basic information
  if( sizeof(FLOAT) == sizeof(double) ){
    b =fscanf(file,"%lf,",&t1); 
    MI->numcomponent=(int)t1;   //number of components 
    b =fscanf(file,"%lf,",&t1); 
    MI->sbin=(int)t1;           //sbin
    b =fscanf(file,"%lf,",&t1);
    MI->interval=(int)t1;       //interval
    b =fscanf(file,"%lf,",&t1); 
    MI->max_Y=(int)t1;          //max_Y
    b =fscanf(file,"%lf,",&t1);				
    MI->max_X=(int)t1;          //max_X
  }else{
    b =fscanf(file,"%f,",&t1);			
    MI->numcomponent=(int)t1;   //number of components 
    b =fscanf(file,"%f,",&t1);		
    MI->sbin=(int)t1;           //sbin
    b =fscanf(file,"%f,",&t1);
    MI->interval=(int)t1;       //interval
    b =fscanf(file,"%f,",&t1); 
    MI->max_Y=(int)t1;          //max_Y
    b =fscanf(file,"%f,",&t1);				
    MI->max_X=(int)t1;          //max_X
  }
  //root filter information
  MI->ridx = (int*)malloc(sizeof(int)*MI->numcomponent);
  MI->oidx = (int*)malloc(sizeof(int)*MI->numcomponent);
  MI->offw = (FLOAT*)malloc(sizeof(FLOAT)*MI->numcomponent);
  MI->rsize = (int*)malloc(sizeof(int)*MI->numcomponent*2);
  MI->numpart = (int*)malloc(sizeof(int)*MI->numcomponent);
  
  //part filter information
  MI->pidx = (int**)malloc(sizeof(int*)*MI->numcomponent);
  MI->didx = (int**)malloc(sizeof(int*)*MI->numcomponent);
  MI->psize = (int**)malloc(sizeof(int*)*MI->numcomponent);
  
  for(int ii=0;ii<MI->numcomponent;ii++)	//LOOP (component)
    {
      if(sizeof(FLOAT)==sizeof(double)) {
        b =fscanf(file,"%lf,",&t1);			
        MI->ridx[ii]=(int)t1-1; //root index
        b =fscanf(file,"%lf,",&t1);			
        MI->oidx[ii]=(int)t1-1; //offset index
        b =fscanf(file,"%lf,",&t1);		
        MI->offw[ii]=t1;        //offset weight (FLOAT)
        b =fscanf(file,"%lf,%lf,",&t1,&t2);
        MI->rsize[ii*2]=(int)t1;   //rsize (Y)
        MI->rsize[ii*2+1]=(int)t2; //rsize (X)
        b =fscanf(file,"%lf,",&t1);		
        MI->numpart[ii]=(int)t1; //number of part filter
      }else{
        b =fscanf(file,"%f,",&t1);			
        MI->ridx[ii]=(int)t1-1; //root index
        b =fscanf(file,"%f,",&t1);			
        MI->oidx[ii]=(int)t1-1; //offset index
        b =fscanf(file,"%f,",&t1);		
        MI->offw[ii]=t1;        //offset weight (FLOAT)
        b =fscanf(file,"%f,%f,",&t1,&t2);
        MI->rsize[ii*2]=(int)t1;   //rsize (Y)
        MI->rsize[ii*2+1]=(int)t2; //rsize (X)
        b =fscanf(file,"%f,",&t1);		
        MI->numpart[ii]=(int)t1; //number of part filter
      }

      MI->pidx[ii]=(int*)malloc(sizeof(int)*MI->numpart[ii]);
      MI->didx[ii]=(int*)malloc(sizeof(int)*MI->numpart[ii]);
      MI->psize[ii]=(int*)malloc(sizeof(int)*MI->numpart[ii]*2);
		
      for(int jj=0;jj<MI->numpart[ii];jj++)	//LOOP (part-filter)
        {
          if(sizeof(FLOAT)==sizeof(double)) {
            b =fscanf(file,"%lf,",&t1);			
            MI->pidx[ii][jj]=(int)t1-1; //part index
            b =fscanf(file,"%lf,",&t1);			
            MI->didx[ii][jj]=(int)t1-1; //define-index of part
            b =fscanf(file,"%lf,%lf,",&t1,&t2);
            MI->psize[ii][jj*2]=(int)t1;
            MI->psize[ii][jj*2+1]=(int)t2;
          }else{
            b =fscanf(file,"%f,",&t1);			
            MI->pidx[ii][jj]=(int)t1-1; //part index
            b =fscanf(file,"%f,",&t1);			
            MI->didx[ii][jj]=(int)t1-1; //define-index of part
            b =fscanf(file,"%f,%f,",&t1,&t2);
            MI->psize[ii][jj*2]=(int)t1;
            MI->psize[ii][jj*2+1]=(int)t2;
          }
        }
    }
  
  //get defs information
  if(sizeof(FLOAT)==sizeof(double)) {
    b =fscanf(file,"%lf,",&t1);	
  }else{
    b =fscanf(file,"%f,",&t1);	
  }
  int DefL = int(t1);
  //MI->def = (FLOAT*)malloc(sizeof(FLOAT)*DefL*4);
  res = cuMemHostAlloc((void **)&(MI->def), sizeof(FLOAT)*DefL*4, CU_MEMHOSTALLOC_DEVICEMAP);
  if(res != CUDA_SUCCESS) {
    printf("cuMemHostAlloc(MI->def) failed: res = %s\n", conv(res));
    exit(1);
  }
  sum_size_def_array = sizeof(FLOAT)*DefL*4;


  MI->anchor = (int*)malloc(sizeof(int)*DefL*2);


  for (int kk=0;kk<DefL;kk++)
    {
      if(sizeof(FLOAT)==sizeof(double)) {
        b =fscanf(file,"%lf,%lf,%lf,%lf,",&t1,&t2,&t3,&t4);	
        MI->def[kk*4]=t1;
        MI->def[kk*4+1]=t2;
        MI->def[kk*4+2]=t3;
        MI->def[kk*4+3]=t4;
        b =fscanf(file,"%lf,%lf,",&t1,&t2);
        MI->anchor[kk*2]=(int)t1;
        MI->anchor[kk*2+1]=(int)t2;
      }else{
        b =fscanf(file,"%f,%f,%f,%f,",&t1,&t2,&t3,&t4);	
        MI->def[kk*4]=t1;
        MI->def[kk*4+1]=t2;
        MI->def[kk*4+2]=t3;
        MI->def[kk*4+3]=t4;
        b =fscanf(file,"%f,%f,",&t1,&t2);
        MI->anchor[kk*2]=(int)t1;
        MI->anchor[kk*2+1]=(int)t2;
      }
    }
  
  //get least_square information
  MI->x1 = (FLOAT **)malloc(sizeof(FLOAT*)*MI->numcomponent);
  MI->x2 = (FLOAT **)malloc(sizeof(FLOAT*)*MI->numcomponent);
  MI->y1 = (FLOAT **)malloc(sizeof(FLOAT*)*MI->numcomponent);
  MI->y2 = (FLOAT **)malloc(sizeof(FLOAT*)*MI->numcomponent);
  
  for(int ii=0;ii<MI->numcomponent;ii++)
    {
      int GL = 1+2*(1+MI->numpart[ii]);
      MI->x1[ii] =(FLOAT *)malloc(sizeof(FLOAT)*GL);
      MI->y1[ii] =(FLOAT *)malloc(sizeof(FLOAT)*GL);
      MI->x2[ii] =(FLOAT *)malloc(sizeof(FLOAT)*GL);
      MI->y2[ii] =(FLOAT *)malloc(sizeof(FLOAT)*GL);
      
      if(sizeof(FLOAT)==sizeof(double)) {
        for (int jj=0;jj<GL;jj++){b =fscanf(file,"%lf,",&t1);	MI->x1[ii][jj]=t1;}
        for (int jj=0;jj<GL;jj++){b =fscanf(file,"%lf,",&t1);	MI->y1[ii][jj]=t1;}
        for (int jj=0;jj<GL;jj++){b =fscanf(file,"%lf,",&t1);	MI->x2[ii][jj]=t1;}
        for (int jj=0;jj<GL;jj++){b =fscanf(file,"%lf,",&t1);	MI->y2[ii][jj]=t1;}
      }else{
        for (int jj=0;jj<GL;jj++){b =fscanf(file,"%f,",&t1);	MI->x1[ii][jj]=t1;}
        for (int jj=0;jj<GL;jj++){b =fscanf(file,"%f,",&t1);	MI->y1[ii][jj]=t1;}
        for (int jj=0;jj<GL;jj++){b =fscanf(file,"%f,",&t1);	MI->x2[ii][jj]=t1;}
        for (int jj=0;jj<GL;jj++){b =fscanf(file,"%f,",&t1);	MI->y2[ii][jj]=t1;}
      }
    }
  
  MI->padx=(int)ceil((double)MI->max_X/2.0+1.0);	//padx
  MI->pady=(int)ceil((double)MI->max_Y/2.0+1.0);	//padY
  
  MI->ini=true;

  
  //fclose
  fclose(file);	
  
  return(MI);
}
Ejemplo n.º 14
0
Partfilters *load_partfilter(char *filename)
{
  FILE *file;		//File
  errno_t err;	//err ( for fopen)
  
  CUresult res;
  
  Partfilters *PF=(Partfilters*)malloc(sizeof(Partfilters));		//Part filter
  
  //fopen	
  //if( (err = fopen_s( &file,filename, "r"))!=0 ) 
  if( (file=fopen(filename, "r"))==NULL ) 
    {
      printf("Part-filter file not found \n");
      exit(-1);
    }
  
  FLOAT t1,t2,t3;
  FLOAT dummy_t1, dummy_t2, dummy_t3;  // variable for dummy scan in order to adjust the location of file-pointer
  
  int b;
  if(sizeof(FLOAT)==sizeof(double)) {
    b =fscanf(file,"%lf,",&t1);		
  }else{
    b =fscanf(file,"%f,",&t1);		
  }
  PF->NoP=(int)t1;											//number of part filter
  
  PF->part_size=(int**)malloc(sizeof(int*)*PF->NoP);			//size of part filter
  PF->partfilter=(FLOAT**)malloc(sizeof(FLOAT*)*PF->NoP);	//weight of part filter
  PF->part_partner=(int*)malloc(sizeof(int)*PF->NoP);			//symmetric information of part
  PF->part_sym=(int*)malloc(sizeof(int)*PF->NoP);				//symmetric information of part
  

  /* keep file pointer location */
  long before_loop_location = ftell(file);

  
  int SUM_SIZE_PART = 0;
  for (int ii=0;ii<PF->NoP;ii++)
    {
      if(sizeof(FLOAT)==sizeof(double)) {
        b =fscanf(file,"%lf,%lf,%lf,",&t1,&t2,&t3);			//number of components 
      }else{
        b =fscanf(file,"%f,%f,%f,",&t1,&t2,&t3);			//number of components 
      }
      PF->part_size[ii]=(int*)malloc(sizeof(int)*3);			
      PF->part_size[ii][0]=(int)t1;
      PF->part_size[ii][1]=(int)t2;
      PF->part_size[ii][2]=(int)t3;
      //printf("***************%f  %f  %f\n",t1,t2,t3);
      int NUMB=PF->part_size[ii][0]*PF->part_size[ii][1]*PF->part_size[ii][2];
#ifdef ORIGINAL
      PF->partfilter[ii]=(FLOAT*)malloc(sizeof(FLOAT)*NUMB);				//weight of root filter
#else
#ifdef SEPARETE_MEM
      res = cuMemHostAlloc((void **)&(PF->partfilter[ii]), sizeof(FLOAT)*NUMB, CU_MEMHOSTALLOC_DEVICEMAP);
      if(res != CUDA_SUCCESS){
        printf("cuMemHostAlloc(PF->partfilter) failed: res = %s\n", conv(res));
        exit(1);
      }
#else
      SUM_SIZE_PART += NUMB*sizeof(FLOAT);
#endif
#endif
     
       /* adjust the location of file-pointer */
      for(int jj=0; jj<NUMB; jj++) {
        if(sizeof(FLOAT)==sizeof(double)) {
          fscanf(file,"%lf,",&dummy_t1);  // this is dummy scan
        }else{
          fscanf(file,"%f,",&dummy_t1);  // this is dummy scan
        }
      }
      if(sizeof(FLOAT)==sizeof(double)) {
        fscanf(file,"%lf,",&dummy_t1); // this is dummy scan
      }else{
        fscanf(file,"%f,",&dummy_t1); // this is dummy scan
      }
      
    }


#ifndef ORIGINAL
#ifndef SEPARETE_MEM
  /* allocate memory region for part in a lump */
  FLOAT *dst_part;
  res = cuMemHostAlloc((void **)&dst_part, SUM_SIZE_PART, CU_MEMHOSTALLOC_DEVICEMAP);
  if(res != CUDA_SUCCESS){
    printf("cuMemHostAlloc(dst_part) failed: res = %s\n", conv(res));
    exit(1);
  }

  /* distribution */
  unsigned long long int pointer = (unsigned long long int)dst_part;
  for(int ii=0; ii<PF->NoP; ii++) {
    PF->partfilter[ii] = (FLOAT *)pointer;
    int NUMB=PF->part_size[ii][0]*PF->part_size[ii][1]*PF->part_size[ii][2];
    pointer += NUMB*sizeof(FLOAT);
  }
#endif
#endif      
      

  /* reset the location of file-pointer */
  fseek(file, before_loop_location, SEEK_SET);
  
  for(int ii=0; ii<PF->NoP; ii++) {  

    int NUMB=PF->part_size[ii][0]*PF->part_size[ii][1]*PF->part_size[ii][2];

    /* adjust the location of file-pointer */
    if(sizeof(FLOAT)==sizeof(double)) {
      fscanf(file,"%lf,%lf,%lf,",&dummy_t1,&dummy_t2,&dummy_t3);  // this is dummy scan
    }else{
      fscanf(file,"%f,%f,%f,",&dummy_t1,&dummy_t2,&dummy_t3);  // this is dummy scan
    }

    for (int jj=0;jj<NUMB;jj++)
      {
        if(sizeof(FLOAT)==sizeof(double)) {
          b =fscanf(file,"%lf,",&t1);		
        }else{
          b =fscanf(file,"%f,",&t1);		
        }
        PF->partfilter[ii][jj]=t1;
      }
    if(sizeof(FLOAT)==sizeof(double)) {
      b =fscanf(file,"%lf,",&t1);		
    }else{
      b =fscanf(file,"%f,",&t1);		
    }
    PF->part_partner[ii]=(int)t1;											//symmetric information of part
    if(PF->part_partner[ii]==0) PF->part_sym[ii]=1;
    else PF->part_sym[ii]=0;
  }
  //fclose
  fclose(file);	
  
  return(PF);
}