Actual source code: ex56.c

petsc-3.14.5 2021-03-03
Report Typos and Errors

  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" ;


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

 33:   PetscInitialize(&argc,&argv,(char*)0,help);if (ierr) return ierr;

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


 41:   VecViennaCLGetCLContext(x, &clctxptr);
 42:   VecViennaCLGetCLQueue(x, &clqueueptr);
 43:   VecViennaCLGetCLMem(x, &clmemptr);

 45:   const cl_context       ctx   = ((const cl_context)clctxptr);
 46:   const cl_command_queue queue = ((const cl_command_queue)clqueueptr);
 47:   const cl_mem           mem   = ((const cl_mem)clmemptr);


 50:   prg = clCreateProgramWithSource(ctx, 1, (const char **) & kernelSrc, NULL, NULL);
 51:   clBuildProgram(prg, 0, NULL, NULL, NULL, NULL);
 52:   knl = clCreateKernel(prg, "doublify", NULL);

 54:   clSetKernelArg(knl, 0, sizeof(cl_mem), &mem);
 55:   clSetKernelArg(knl, 1, sizeof(PetscInt), &size);

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

 61:   // let petsc know that device data is altered
 62:   VecViennaCLRestoreCLMem(x);

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

 67:   VecDestroy(&x);
 68:   clReleaseContext(ctx);
 69:   clReleaseCommandQueue(queue);
 70:   clReleaseMemObject(mem);
 71:   clReleaseProgram(prg);
 72:   clReleaseKernel(knl);

 74:   PetscFinalize();
 75:   return ierr;
 76: }

 78: /*TEST

 80:    build:
 81:       requires: viennacl

 83:    test:
 84:       nsize: 1
 85:       suffix: 1
 86:       args: -viennacl_backend opencl -viennacl_opencl_device_type gpu

 88: TEST*/