title: OpenCL problem with double type
body: <p>I'm trying use double type in openCL, but doesn't work anyway, i want use double for more precision, if have any other type make this, please, tell me.</p>
  
  <p>if you don't have time for read my code, resuming is: I want use double(or other type) in openCL for more precision in calculation of pi.</p>
  
  <p>My code:</p>
  
  <pre><code> #pragma OPENCL EXTENSION cl_amd_fp64 : enable
  
   #include &lt;fcntl.h&gt;
   #include &lt;stdio.h&gt;
   #include &lt;stdlib.h&gt;
   #include &lt;string.h&gt;
   #include &lt;math.h&gt;
   #include &lt;unistd.h&gt;
   #include &lt;sys/types.h&gt;
   #include &lt;sys/stat.h&gt;
   #include &lt;OpenCL/opencl.h&gt;
  
   ////////////////////////////////////////////////////////////////////////////////
  
   // Use a static data size for simplicity
   //
   #define DATA_SIZE (1000000)
   ////////////////////////////////////////////////////////////////////////////////
   #define TIPO double
   // Simple compute kernel that computes the calcpi of an input array. [1]
   //
   const char *KernelSource = "\n" \
   "#pragma OPENCL EXTENSION cl_amd_fp64 : enable \n" \
   "__kernel void calcpi( \n" \
   " __global double* input, \n" \
   " __global double* output, \n" \
   " const unsigned int count) \n" \
   "{ \n" \
   " int i = get_global_id(0); \n" \
   " double z = get_global_id(0)*2+1; \n" \
   " if(i &lt; count) \n" \
   " output[i] = 4.0/z; \n" \
   "} \n" \
   "\n";
  
   ////////////////////////////////////////////////////////////////////////////////
  
   int main(int argc, char** argv)
   {
   int err; // error code returned from api calls
   //printf("%d",sizeof(TIPO));
   //scanf("%d",&amp;err);
   TIPO data[2]; // original data set given to device
   TIPO *results = malloc(sizeof(TIPO)*DATA_SIZE); // results returned from device
   //unsigned int correct; // number of correct results returned
   //printf("TESTE");
  
   size_t global; // global domain size for our calculation
   size_t local; // local domain size for our calculation
  
   cl_device_id device_id; // device ID
   cl_context context; // context
   cl_command_queue queue; // command queue
   cl_program program; // program
   cl_kernel kernel; // kernel
  
   cl_mem input; // device memory used for the input array
   cl_mem output; // device memory used for the output array
  
   // Get data on which to operate
   //
  
   //int i = 0;
   //int n = 3;
   unsigned int count = DATA_SIZE;
   //for(i = 0; i &lt; count; i+=2) {
   //data[i] = n;
   //n += 2;
   //}
   //printf("TESTE");
   // Get an ID for the device [2]
   int gpu = 1;
   err = clGetDeviceIDs(NULL, gpu ? CL_DEVICE_TYPE_GPU : CL_DEVICE_TYPE_CPU, 1,&amp;device_id,      NULL);
   if (err != CL_SUCCESS)
        printf("ERROR CLGETDEVICEIDS!\n");     // [3]
  
   // Create a context [4]
   //
   context = clCreateContext(0, 1, &amp;device_id, NULL, NULL, &amp;err);
   if (!context) {
        printf("ERROR CONTEXT\n");
   }
  
   // Create a command queue [5]
   //
   queue = clCreateCommandQueue(context, device_id, 0, &amp;err);
   if (!queue) {
        printf("ERROR QUEUE\n");
   }
  
   // Create the compute program from the source buffer [6]
   //
   program = clCreateProgramWithSource(context, 1,(const char **) &amp; KernelSource, NULL, &amp;err);
   if ( !program) {
        printf("ERROR PROGRAM\n");
   }
  
   // Build the program executable [7]
   //
   err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
   if (err != CL_SUCCESS)
   {
        size_t len;
        char buffer[2048];
  
        printf("Error: Failed to build program executable\n"); //[8]
        clGetProgramBuildInfo(program, device_id, CL_PROGRAM_BUILD_LOG,sizeof(buffer), buffer, &amp;len);
        printf("%s\n", buffer);
        exit(1);
   }
  
   // Create the compute kernel in the program we wish to run [9]
   //
   kernel = clCreateKernel(program, "calcpi", &amp;err);
   if (!kernel || err != CL_SUCCESS) {
        printf("ERROR KERNEL OR CL_SUCESS\n");
   }
  
   // Create the input and output arrays in device memory for our calculation
   // [10]
   input = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(TIPO) *count,NULL, NULL);
   output = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(TIPO) *count,NULL, NULL);
   if (!input || !output) {
        printf("ERROR !INPUT OR !OUTPUT\n");
   }
  
   // Write our data set into the input array in device memory [11]
   //
   err = clEnqueueWriteBuffer(queue, input, CL_TRUE, 0,sizeof(TIPO) *2, data, 0, NULL, NULL);
   if (err != CL_SUCCESS) {
        printf("ERROR WRITE OUR DATA\n");
   }
  
   // Set the arguments to our compute kernel [12]
   //
   err = 0;
   err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &amp;input);
   err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), &amp;output);
   err |= clSetKernelArg(kernel, 2, sizeof(unsigned int), &amp;count);
   if (err != CL_SUCCESS) {
        printf("ERROR ARGUMENTS COMPUTE KERNEL - ERROR NUMBER: %d\n",err);
        exit(1);
   }
  
   // Get the maximum work-group size for executing the kernel on the device
   // [13]
   err = clGetKernelWorkGroupInfo(kernel, device_id, CL_KERNEL_WORK_GROUP_SIZE,sizeof(size_t), &amp;local, NULL);
   if (err != CL_SUCCESS) {
        printf("ERROR MAXIMUM WORK-GROUP - ERROR NUMBER: %d\n",err);
        exit(1);
   }
  
   // Execute the kernel over the entire range of the data set [14]
   //
   global = count;
   //printf("TESTE");
   err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &amp;global, NULL,0, NULL, NULL);
   if (err) {
        printf("ERROR EXECUTE KERNEL - ERROR NUMBER: %d\n",err);
        exit(1);
   }
  
   // Wait for the command queue to get serviced before reading back results
   // [15]
   clFinish(queue);
  
   // Read the results from the device [16]
   //
   err = clEnqueueReadBuffer(queue, output, CL_TRUE, 0,sizeof(TIPO) *count, results, 0, NULL, NULL );
   if (err != CL_SUCCESS) {
        printf("ERROR READ RESULTS - ERROR NUMBER: %d\n",err);
   }
   //printf("TESTE");
   TIPO pi = 0.0;
   int i;
   for (i=0;i&lt;count-1;i++) {
        //printf("%f",results[i]);
        pi += (pow(-1.0,i)) * (TIPO) results[i];
        //pi = (TIPO) results[i];
        //printf("casa %d deu: %1.50f\n",i,pi);
        //printf("%f",(pow(-1,i)));
        //pi += (pow(-1.0,i));
   }
   printf("PI: %1.50f",pi);
  
   // Shut down and clean up
   //
   clReleaseMemObject(input);
   clReleaseMemObject(output);
   clReleaseProgram(program);
   clReleaseKernel(kernel);
   clReleaseCommandQueue(queue);
   clReleaseContext(context);
  
   scanf("%d",&amp;i);
   return 0;
   }
  </code></pre>
  
  <p>when i put in kernelSource:</p>
  
  <p>output = 4.0;</p>
  
  <p>only like this, i get 512.000123023986816406250000000000000000000000000 </p>
  
  <p>in results..</p>
  
  <p>or 1.0 = 0.00781250184809323400259017944335937500000000000</p>