Actual source code: ex56.c


  2: static char help[] = "Update the data in a VECVIENNACL via a CL kernel.\n\n";

  4: #include <petscvec.h>
  5: #include <CL/cl.h>

  7: const char *kernelSrc =                                       "\n" \
  8: "#pragma OPENCL EXTENSION cl_khr_fp64 : enable                 \n" \
  9: "__kernel void doublify(  __global double *x,                  \n" \
 10: "                       const unsigned int n)                  \n" \
 11: "{                                                             \n" \
 12: "  //Get our global thread ID                                  \n" \
 13: "  int gid = get_global_id(0);                                 \n" \
 14: "                                                              \n" \
 15: "  if (gid < n)                                                \n" \
 16: "    x[gid] = 2*x[gid];                                        \n" \
 17: "}                                                             \n" \
 18:                                                               "\n" ;

 20: int main(int argc,char **argv)
 21: {
 22:   PetscInt          size=5;
 23:   Vec               x;
 24:   cl_program        prg;
 25:   cl_kernel         knl;
 26:   PETSC_UINTPTR_T   clctxptr;
 27:   PETSC_UINTPTR_T   clqueueptr;
 28:   PETSC_UINTPTR_T   clmemptr;
 29:   const size_t      gsize=10, lsize=2;

 31:   PetscInitialize(&argc,&argv,(char*)0,help);

 33:   VecCreate(PETSC_COMM_WORLD,&x);
 34:   VecSetSizes(x,size,PETSC_DECIDE);
 35:   VecSetType(x, VECVIENNACL);
 36:   VecSet(x, 42.0);

 38:   VecViennaCLGetCLContext(x, &clctxptr);
 39:   VecViennaCLGetCLQueue(x, &clqueueptr);
 40:   VecViennaCLGetCLMem(x, &clmemptr);

 42:   const cl_context       ctx   = ((const cl_context)clctxptr);
 43:   const cl_command_queue queue = ((const cl_command_queue)clqueueptr);
 44:   const cl_mem           mem   = ((const cl_mem)clmemptr);

 46:   prg = clCreateProgramWithSource(ctx, 1, (const char **) & kernelSrc, NULL, NULL);
 47:   clBuildProgram(prg, 0, NULL, NULL, NULL, NULL);
 48:   knl = clCreateKernel(prg, "doublify", NULL);

 50:   clSetKernelArg(knl, 0, sizeof(cl_mem), &mem);
 51:   clSetKernelArg(knl, 1, sizeof(PetscInt), &size);

 53:   // Launch the kernel. (gsize > size: masked execution of some work items)
 54:   clEnqueueNDRangeKernel(queue, knl, 1, NULL, &gsize, &lsize, 0, NULL, NULL);
 55:   clFinish(queue);

 57:   // let petsc know that device data is altered
 58:   VecViennaCLRestoreCLMem(x);

 60:   // 'x' should contain 84 as all its entries
 61:   VecView(x,PETSC_VIEWER_STDOUT_WORLD);

 63:   VecDestroy(&x);
 64:   clReleaseContext(ctx);
 65:   clReleaseCommandQueue(queue);
 66:   clReleaseMemObject(mem);
 67:   clReleaseProgram(prg);
 68:   clReleaseKernel(knl);

 70:   PetscFinalize();
 71:   return 0;
 72: }

 74: /*TEST

 76:    build:
 77:       requires: viennacl

 79:    test:
 80:       nsize: 1
 81:       suffix: 1
 82:       args: -viennacl_backend opencl -viennacl_opencl_device_type gpu

 84: TEST*/