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:   PetscFunctionBeginUser;
 32:   PetscCall(PetscInitialize(&argc, &argv, (char *)0, help));

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

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

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

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

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

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

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

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

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

 71:   PetscCall(PetscFinalize());
 72:   return 0;
 73: }

 75: /*TEST

 77:    build:
 78:       requires: viennacl

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

 85: TEST*/