// create a compute context with GPU devicecontext = clCreateContextFromType(CL_DEVICE_TYPE_GPU);// create a work-queuequeue = clCreateWorkQueue(context, NULL, NULL, 0);// allocate the buffer memory objectsmemobjs[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeof(float)*2*num_entries, srcA);memobjs[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(float)*2*num_entries, NULL);// create the compute programprogram = clCreateProgramFromSource(context, 1, &fft1D_1024_kernel_src, NULL);// build the compute program executableclBuildProgramExecutable(program, false, NULL, NULL);// create the compute kernelkernel = clCreateKernel(program, "fft1D_1024");// create N-D range object with work-item dimensionsglobal_work_size[0] = n;local_work_size[0] = 64;range = clCreateNDRangeContainer(context, 0, 1, global_work_size, local_work_size);// set the args valuesclSetKernelArg(kernel, 0, (void *)&memobjs[0], sizeof(cl_mem), NULL);clSetKernelArg(kernel, 1, (void *)&memobjs[1], sizeof(cl_mem), NULL);clSetKernelArg(kernel, 2, NULL, sizeof(float)*(local_work_size[0]+1)*16, NULL);clSetKernelArg(kernel, 3, NULL, sizeof(float)*(local_work_size[0]+1)*16, NULL); // execute kernelclExecuteKernel(queue, kernel, NULL, range, NULL, 0, NULL);

// This kernel computes FFT of length 1024. The 1024 length FFT is decomposed into// calls to a radix 16 function, another radix 16 function and then a radix 4 function__kernel void fft1D_1024 (__global float2 *in, __global float2 *out,                          __local float *sMemx, __local float *sMemy) {  int tid = get_local_id(0);  int blockIdx = get_group_id(0) * 1024 + tid;  float2 data[16];  // starting index of data to/from global memory  in = in + blockIdx;  out = out + blockIdx;  globalLoads(data, in, 64); // coalesced global reads  fftRadix16Pass(data);      // in-place radix-16 pass  twiddleFactorMul(data, tid, 1024, 0);  // local shuffle using local memory  localShuffle(data, sMemx, sMemy, tid, (((tid & 15) * 65) + (tid >> 4)));  fftRadix16Pass(data);               // in-place radix-16 pass  twiddleFactorMul(data, tid, 64, 4); // twiddle factor multiplication  localShuffle(data, sMemx, sMemy, tid, (((tid >> 4) * 64) + (tid & 15)));  // four radix-4 function calls  fftRadix4Pass(data); fftRadix4Pass(data + 4);  fftRadix4Pass(data + 8); fftRadix4Pass(data + 12);  // coalesced global writes  globalStores(data, out, 64);}