Actual source code: ex56.c

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

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

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

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

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

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

 38:   PetscCall(VecViennaCLGetCLContext(x, &clctxptr));
 39:   PetscCall(VecViennaCLGetCLQueue(x, &clqueueptr));
 40:   PetscCall(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:   PetscCall(VecViennaCLRestoreCLMem(x));

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

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

 70:   PetscCall(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*/