Esempio n. 1
0
int
Lucas::fft_setup (int length)
{
  clAmdFftSetupData fftSetupData;
  OPENCL_V_THROW(clAmdFftInitSetupData (&fftSetupData),"Failed to clAmdFftInitSetupData.");
  fftSetupData.debugFlags = CLFFT_DUMP_PROGRAMS;        // Dumps the FFT kernels
  // Setup the AMD FFT library.
  OPENCL_V_THROW(clAmdFftSetup (&fftSetupData),"Failed to clAmdFftSetup.");
  // Create FFT Plan
  const size_t logicalDimensions[1] = { length };
  // Create default plan.
  OPENCL_V_THROW(clAmdFftCreateDefaultPlan(&plan, context(), CLFFT_1D, logicalDimensions),"Failed to clAmdFftCreateDefaultPlan.");
  // Set double precision.
  OPENCL_V_THROW(clAmdFftSetPlanPrecision(plan, CLFFT_DOUBLE),"Failed to clAmdFftSetPlanPrecision.");
  // Set layout.
  OPENCL_V_THROW(clAmdFftSetLayout(plan, CLFFT_COMPLEX_INTERLEAVED,CLFFT_COMPLEX_INTERLEAVED),"Failed to clAmdFftSetLayout.");
  // Normalize forward transformation.
  OPENCL_V_THROW(clAmdFftSetPlanScale(plan, CLFFT_FORWARD, 1.0f/static_cast<cl_float>(length)),"Failed to clAmdFftSetPlanScale.");
  // Normalize backward transformation.
  OPENCL_V_THROW(clAmdFftSetPlanScale(plan, CLFFT_BACKWARD, 1.0f),"Failed to clAmdFftSetPlanScale.");
  // In-place FFT.
  OPENCL_V_THROW(clAmdFftSetResultLocation(plan, CLFFT_INPLACE),"Failed to clAmdFftSetResultLocation.");
  // Set number of transformations per plan.
  OPENCL_V_THROW(clAmdFftSetPlanBatchSize(plan, 1),"Failed to clAmdFftSetPlanBatchSize.");
  // BakePlan
  OPENCL_V_THROW(clAmdFftBakePlan(plan, 1, &commandQueue(), NULL, NULL),"Failed to clAmdFftBakePlan.");
}
int main()
{
        //Control Variables
        bool showStartInput=false;// Setting it to true shows the original Input
        bool showFftOutput=false;// Shows the output after the FFT but before the Reshuffle
        bool showReshuffleOutput=false;// Shows the output after the reshuffle
        bool showFinalResult=false; // Shows final result after cross-correlation
        bool showGemmInput=false; // Shows output after the reshuffle but before the matrix multiplication
	bool showReformatOutputAfterReshuffle=false; // Shows output after it has been reformatted after the reshuffling

	//openCL State
        cl_platform_id platform_id=NULL;
	cl_device_id device_id=NULL;
        cl_context context=NULL;
        cl_command_queue queue=NULL;
        cl_program program=NULL;
        cl_kernel kernel=NULL;
        cl_uint ret_num_devices;
        cl_uint ret_num_platforms;
        cl_int ret=0; // Stores the error values retuned by many functions
	cl_event event = NULL;
        cl_event events[10];
	cl_kernel clKernel;

	//FFT state
	clAmdFftPlanHandle plHandle;
        clAmdFftResultLocation place = CLFFT_OUTOFPLACE; //Alternative CLFFT_INPLACE
        clAmdFftLayout inLayout  = CLFFT_COMPLEX_INTERLEAVED;
        clAmdFftLayout outLayout = CLFFT_COMPLEX_INTERLEAVED;
        clAmdFftDim dim = CLFFT_1D;
	size_t clStrides[3]={0,0,0};
        size_t clLengths[3];
	clLengths[0]=(MEM_SIZE/2);//Length of first dimension of fft
	clLengths[1]=1;//length of second dimension of fft
	clLengths[2]=1;
	clStrides[ 0 ] = 1;
        clStrides[ 1 ] = clStrides[ 0 ] * clLengths[ 0 ];
        clStrides[ 2 ] = clStrides[ 1 ] * clLengths[ 1 ];
        clStrides[ 3 ] = clStrides[ 2 ] * clLengths[ 2 ];
	size_t batchSize=CHANSIZE;//number of discreet fft's to be calculated simultaneously


	//Initialise openCL
	OPENCL_V_THROW(clGetPlatformIDs(1, &platform_id, &ret_num_platforms),"clGetPlatformIDs Failed");
        OPENCL_V_THROW(clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id,&ret_num_devices),"clGetDeviceIDs Failed");
        context = clCreateContext(NULL, 1, &device_id, NULL, NULL, &ret);
        OPENCL_V_THROW(ret, "Creating Context failed" );
        queue = clCreateCommandQueue(context, device_id, 0, &ret);
        OPENCL_V_THROW(ret, "Creating command queue failed" );


//===========Initialise the host buffers======================================

/*
 * The functions sgenerate2darray(), screate2darray() and sgenerate2darrayout() are defined and declared in definition.h
 */

	float** src_a_h=sgenerate2darray(NO_INPUTS,MEM_SIZE);//To be used to store the original input

	float** answer=screate2darray(NO_INPUTS,MEM_SIZE);//To be used to store the answer after the reshuffling

        float** corr_h=sgenerate2darrayout(NO_INPUTS,CHANSIZE << 1,CHANNELNO);// To be used to store the final answer

        if(showStartInput){
            cout << "Initial Input Buffer" << "\n";
            for(int j=0;j<NO_INPUTS;j++){
                for(int i=0;i<MEM_SIZE;i++){
                        cout << src_a_h[j][i] << " ";
                }cout << "\n";
            }printf("\n");
        }


//===================================================================
        //Calculation of facs for reshuffling
	complex <float>* facs_h=(complex <float>*) malloc(sizeof(complex <float>)*(MEM_SIZE/2));
	complex<float> I=1.0i;
        complex <float> xx=2.0*PI;
        for(int i=0;i<MEM_SIZE/2;i++){
                facs_h[i]=(1.0*i)/(1.0*MEM_SIZE);
                facs_h[i]=exp(xx*(-I*facs_h[i]));
        }
//===================================================================
	//Initialise GPU memory buffers
	size_t sizeofgpumem=NO_INPUTS*MEM_SIZE*sizeof(float);
	size_t sizeoffacsmem=MEM_SIZE*sizeof(float);
	cl_mem clMemBuffersIn = clCreateBuffer(context,CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,sizeofgpumem,src_a_h[0],&ret);
	OPENCL_V_THROW( ret, "Creating clMemBuffersIn Buffer failed" );

	cl_mem clMemBuffersOut = clCreateBuffer(context,CL_MEM_READ_WRITE,sizeofgpumem,NULL,&ret);
	OPENCL_V_THROW (ret, "Creating fft output Buffer failed");

	cl_mem facs = clCreateBuffer(context,CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,sizeoffacsmem,facs_h,&ret);
	OPENCL_V_THROW (ret, "Creating facs Buffer failed");

	//===========================Starting the fft=============================//
	clAmdFftSetupData setupData;
	OPENCL_V_THROW( clAmdFftInitSetupData( &setupData ),"clAmdFftInitSetupData failed" );
        OPENCL_V_THROW( clAmdFftSetup( &setupData ), "clAmdFftSetup failed" );
	OPENCL_V_THROW( clAmdFftCreateDefaultPlan( &plHandle, context, dim, clLengths ), "clAmdFftCreateDefaultPlan failed" );
	OPENCL_V_THROW (clAmdFftSetPlanBatchSize (plHandle, batchSize),"Setting BatchSize Failed");
        OPENCL_V_THROW (clAmdFftSetResultLocation( plHandle, place ), "clAmdFftSetResultLocation failed" );
	OPENCL_V_THROW (clAmdFftSetPlanInStride  ( plHandle, dim, clStrides ), "clAmdFftSetPlanInStride failed" );
        OPENCL_V_THROW (clAmdFftSetPlanOutStride ( plHandle, dim, clStrides ), "clAmdFftSetPlanOutStride failed" );
        OPENCL_V_THROW (clAmdFftSetPlanDistance  ( plHandle, clStrides[ dim ], clStrides[ dim ]), "clAmdFftSetPlanDistance failed" );

	OPENCL_V_THROW( clAmdFftBakePlan( plHandle, 1, &queue, NULL, NULL ), "clAmdFftBakePlan failed" );

	size_t tempbuffersize=0;
        OPENCL_V_THROW( clAmdFftGetTmpBufSize(plHandle, &tempbuffersize ), "clAmdFftGetTmpBufSize failed" );

	//allocate the intermediate buffer
        cl_mem clMedBuffer=NULL;
        if (tempbuffersize)
        {
                cl_int medstatus;
                clMedBuffer = clCreateBuffer ( context, CL_MEM_READ_WRITE | CL_MEM_ALLOC_HOST_PTR,tempbuffersize, 0, &medstatus);
                OPENCL_V_THROW( medstatus, "Creating fft intermediate Buffer failed" );
        }
	if (( place == CLFFT_INPLACE )&& ( inLayout != outLayout )) {
                switch( inLayout )
                {
                case CLFFT_COMPLEX_INTERLEAVED:
                        {
                                assert (CLFFT_COMPLEX_PLANAR == outLayout);
                                throw std::runtime_error( "Cannot use the same buffer for interleaved->planar in-place transforms" );
                                break;
                        }
                case CLFFT_COMPLEX_PLANAR:
                        {
                                assert (CLFFT_COMPLEX_INTERLEAVED == outLayout);
                                throw std::runtime_error( "Cannot use the same buffer for planar->interleaved in-place transforms" );
                                break;
                        }
                }
        }
	cl_mem * BuffersOut = ( place == CLFFT_INPLACE ) ? NULL :  &clMemBuffersOut;

	//========Timimg fft============//

	double time_fft_start=omp_get_wtime();

	for(int i=0;i<ITER_FFT;i++){
		OPENCL_V_THROW( clAmdFftEnqueueTransform( plHandle, CLFFT_FORWARD, 1,&queue,0,NULL,&event,&clMemBuffersIn,BuffersOut,clMedBuffer ),"clAmdFftEnqueueTransform failed" );
	}
	ret=clWaitForEvents(1,&event);

        double time_fft_end=omp_get_wtime();

	//Cleaning up fft
        OPENCL_V_THROW( clAmdFftDestroyPlan( &plHandle ), "clAmdFftDestroyPlan failed" );
        OPENCL_V_THROW( clAmdFftTeardown( ), "clAmdFftTeardown failed" );

        //displaying results
        if(showFftOutput){
            OPENCL_V_THROW( clEnqueueReadBuffer( queue, clMemBuffersOut, CL_TRUE, 0, sizeofgpumem,answer [0], 0, NULL, NULL ),"Reading the result buffer failed" );

            cout << "**FFT Output**" << endl;
            for(int j=0;j<NO_INPUTS;j++){
                    for(int i=0;i<MEM_SIZE;i++){
                            cout << answer[j][i] << " ";
                    } printf("\n");
            }printf("\n");
        }
	//==================End of FFT=============================================//

	//==================Start the Reshuffling==================================//
	FILE *fp;
	char fileName[]="./reshuffle.cl";
	char* source_str=NULL;
	size_t source_size;

	//Load the source code containing the kernel/
 	fp = fopen(fileName, "r");
 	if (!fp) {
 		fprintf(stderr, "Failed to load reshuffle kernel.¥n");
	 	exit(1);
 	}
	source_str = (char*)malloc(MAX_SOURCE_SIZE);
	source_size = fread(source_str, 1, MAX_SOURCE_SIZE, fp);
	fclose(fp);

	//Preparation for building the Kernel
	program = clCreateProgramWithSource(context, 1, (const char **)&source_str,(const size_t *)&source_size, &ret);
        OPENCL_V_THROW( ret, "Creating program with source failed for Reshuffle" );
	OPENCL_V_THROW( clBuildProgram(program, 1, &device_id, NULL, NULL, NULL),"Build Program Failed for Reshuffle");
	kernel = clCreateKernel(program, "reshuffle", &ret);
	OPENCL_V_THROW( ret, "Creating kernel failed for Reshuffle" );

	//Set kernel parameters
	const int num=NO_INPUTS*MEM_SIZE;
	const int block=MEM_SIZE;
	OPENCL_V_THROW(clSetKernelArg(kernel, 0, sizeof(cl_mem), (float *)&clMemBuffersIn),"Passing argument 0 of reshuffle failed");
 	OPENCL_V_THROW(clSetKernelArg(kernel, 1, sizeof(cl_mem), (float *)&facs),"Passing arg 1 of reshuffle failed");
        OPENCL_V_THROW(clSetKernelArg(kernel, 2, sizeof(cl_mem), (float *)&clMemBuffersOut),"Passing arg2 of reshuffle failed");
 	OPENCL_V_THROW(clSetKernelArg(kernel, 3, sizeof(int), (int *)&num),"Passing arg3 of reshuffle failed");
        OPENCL_V_THROW(clSetKernelArg(kernel, 4, sizeof(int), (int *)&block),"Passing arg4 of reshuffle failed");


	// Execute OpenCL Kernel //
	const size_t local_ws=NO_INPUTS*MEM_SIZE;
	const size_t global_ws=min(NO_THREAD_PER_BLOCK,MEM_SIZE);//ceil(MEM_SIZE/local_ws);

	//===========timing the reshuffle===============//
	double time_reshuffle_start=omp_get_wtime();
	for(int i=0;i<ITER_FFT;i++){
		OPENCL_V_THROW(clEnqueueNDRangeKernel(queue,kernel, 1, NULL,&local_ws,&global_ws, 0, NULL, NULL),"Reshuffle Kernel execution failed");
	}
	double time_reshuffle_end=omp_get_wtime();

	//Read back data
	OPENCL_V_THROW(clEnqueueReadBuffer(queue, clMemBuffersOut, CL_TRUE, 0, sizeofgpumem,answer[0], 0, NULL, NULL),"Reading back reshuffled data failed");

	//====================Finish the reshuffling================================//

        if(showReshuffleOutput){
                cout << "Output after reshuffling" << endl;
                for(int j=0;j<NO_INPUTS;j++){
                        for(int i=0;i<MEM_SIZE;i++){
                                cout << answer[j][i] << " ";
                        } printf("\n");
                }printf("\n");
        }
        //=================Reformatting the input given to the matrix multiply===================================//
        float** answer_final=screate2darray(NO_INPUTS*2,MEM_SIZE/2);
        for(int i=0;i<NO_INPUTS;i++){
            for(int j=0;j<MEM_SIZE;j++){
                if(j&1)
                        answer_final[(i<<1)+1][j >> 1]=answer[i][j];
                else
                        answer_final[(i<<1)][j >> 1]=answer[i][j];
            }
        }
Esempio n. 3
0
clAmdFftPlanHandle CreateFFTPlan(clLabviewDevice *d, 
					int FFTType, int Dimension,  
					int Width, int Height, int Depth,
					int StrideW, int StrideH, int StrideD, int StrideT, 
					int PaddingW, int PaddingH, int PaddingD, 
					size_t *OutputWidthFloat, int SingleOrDouble, int *Error){
	
	clAmdFftPlanHandle	plHandle = 0;

#ifndef NO_OPENCL
	*Error = clLabviewDevice::Error(clLabviewDevice::SanitizeDevice(d));
	if(*Error != 0)
		return NULL;					
						
	clAmdFftResultLocation	place = CLFFT_OUTOFPLACE;       //INPLACE NOT SUPPORTED
	size_t clStridesIn[ 4 ];
	size_t clStridesOut[ 4 ];
	clAmdFftLayout inLayout;
	clAmdFftLayout outLayout; 
	size_t batchSize;
	clAmdFftDim	dim;
	size_t clLengths[ 3 ];

	clAmdFftStatus status;

	int DistanceIn, DistanceOut;
	int OutputWidth;

	int FFTSize;

	if(FFTType == 1){
		//Needed to the FFT does the proper size on the inverse
		FFTSize = (Width - 1)*2;
	}else{
		FFTSize = Width;
	}

	switch(Dimension){
		case 0:
			//1D FFT
			batchSize = Depth*Height;
			dim = CLFFT_1D;
			clLengths[0] = FFTSize;
			clLengths[1] = 1;
			clLengths[2] = 1;
		break;

		case 1:
			//2D FFT
			batchSize = Depth;	
			dim = CLFFT_2D;
			clLengths[0] = FFTSize;
			clLengths[1] = Height;
			clLengths[2] = 1;
		break;
	
		case 2:
			//3D FFT
			batchSize = 1;
			dim = CLFFT_3D;
			clLengths[0] = FFTSize;
			clLengths[1] = Height;
			clLengths[2] = Depth;
		break;
	}

	switch(FFTType){
		//These cases are defined by the labview Type Def
		
		case 0:
			inLayout = CLFFT_REAL;
			outLayout = CLFFT_HERMITIAN_PLANAR;
			//*BufferSize = clLengths[2]*clLengths[1]*(clLengths[0]/2 + 1)*batchSize*sizeof(std::complex<float>)/sizeof(float);
			OutputWidth = FFTSize/2 + 1;
			*OutputWidthFloat = OutputWidth;
			DistanceIn = Width*clLengths[1]*clLengths[2];
			DistanceOut = OutputWidth*clLengths[1]*clLengths[2];
		break;

		case 1:
			inLayout = CLFFT_HERMITIAN_PLANAR;
			outLayout = CLFFT_REAL;
			//*BufferSize = clLengths[2]*clLengths[1]*(clLengths[0] - 1)*2*batchSize;
			OutputWidth = FFTSize;
			*OutputWidthFloat = OutputWidth;
			DistanceIn = Width*clLengths[1]*clLengths[2];
			DistanceOut = OutputWidth*clLengths[1]*clLengths[2];
		break;			
		
		case 2:
			inLayout = CLFFT_COMPLEX_PLANAR;
			outLayout = CLFFT_COMPLEX_PLANAR;
			//*BufferSize = clLengths[2]*clLengths[1]*(clLengths[0])*batchSize*sizeof(std::complex<float>)/sizeof(float);
			OutputWidth = FFTSize;
			*OutputWidthFloat = OutputWidth;
			DistanceIn = Width*clLengths[1]*clLengths[2];
			DistanceOut = OutputWidth*clLengths[1]*clLengths[2];
			place = CLFFT_INPLACE;
		break;
	}

	size_t fftVectorSize= 0, fftVectorSizePadded = 0, fftBatchSize = 0;
	
	clStridesIn[0] = 1;
	clStridesIn[ 1 ] = (clStridesIn[ 0 ])*(Width + PaddingW);
	clStridesIn[ 2 ] = (clStridesIn[ 1 ])*(clLengths[ 1 ] + PaddingH);
	clStridesIn[ 3 ] = (clStridesIn[ 2 ])*(clLengths[ 2 ] + PaddingD);

	clStridesOut[0] = 1;
	clStridesOut[ 1 ] = (clStridesOut[ 0 ])*(OutputWidth + PaddingW);
	clStridesOut[ 2 ] = (clStridesOut[ 1 ])*(clLengths[ 1 ] + PaddingH);
	clStridesOut[ 3 ] = (clStridesOut[ 2 ])*(clLengths[ 2 ] + PaddingD);

	fftVectorSize	= clLengths[ 0 ] * clLengths[ 1 ] * clLengths[ 2 ];
	fftVectorSizePadded = clStridesIn[ 3 ];
	fftBatchSize	= fftVectorSizePadded * batchSize;

	status = clAmdFftCreateDefaultPlan( &plHandle, d->GetContext(), dim, clLengths );

	if(status != 0){
		*Error = clLabviewDevice::Error(OPENCLV_FFT_DEFAULT_PLAN_FAILED);
		return NULL;
	}

	status = clAmdFftSetResultLocation( plHandle, place );

	if(status != 0){
		*Error = clLabviewDevice::Error(OPENCLV_FFT_SET_RESULT_FAILED);
		return NULL;
	}

	status = clAmdFftSetLayout( plHandle, inLayout, outLayout );

	if(status != 0){
		*Error = clLabviewDevice::Error(OPENCLV_FFT_SET_LAYOUT_FAILED);
		return NULL;
	}

	if(SingleOrDouble == 0){
		status = clAmdFftSetPlanPrecision(plHandle, CLFFT_SINGLE);
	}else{
		status = clAmdFftSetPlanPrecision(plHandle, CLFFT_DOUBLE);
	}

	if(status != 0){
		*Error = clLabviewDevice::Error(OPENCLV_FFT_SET_PRECISION_FAILED);
		return NULL;
	}

	status = clAmdFftSetPlanBatchSize( plHandle, batchSize );

	if(status != 0){
		*Error = clLabviewDevice::Error(OPENCLV_FFT_SET_BATCHSIZE_FAILED);
		return NULL;
	}
	
	status = clAmdFftSetPlanInStride  ( plHandle, dim, clStridesIn );

	if(status != 0){
		*Error = clLabviewDevice::Error(OPENCLV_FFT_SET_STRIDES_IN_FAILED);
		return NULL;
	}

	status = clAmdFftSetPlanOutStride ( plHandle, dim, clStridesOut );

	if(status != 0){
		*Error = clLabviewDevice::Error(OPENCLV_FFT_SET_STRIDES_OUT_FAILED);
		return NULL;
	}

	status = clAmdFftSetPlanDistance  ( plHandle, DistanceIn, DistanceOut);

	if(status != 0){
		*Error = clLabviewDevice::Error(OPENCLV_FFT_SET_PLAN_DIST_FAILED);
		return NULL;
	}

	status = clAmdFftBakePlan( plHandle, 0, d->GetQueuePtr(), NULL, NULL );

	if(status != 0){
		*Error = clLabviewDevice::Error(OPENCLV_FFT_BAKE_PLAN_FAILED);
		return NULL;
	}

#endif
	
	return plHandle;


}