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;
32: PetscInitialize(&argc, &argv, (char *)0, help);
34: VecCreate(PETSC_COMM_WORLD, &x);
35: VecSetSizes(x, size, PETSC_DECIDE);
36: VecSetType(x, VECVIENNACL);
37: VecSet(x, 42.0);
39: VecViennaCLGetCLContext(x, &clctxptr);
40: VecViennaCLGetCLQueue(x, &clqueueptr);
41: 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: VecViennaCLRestoreCLMem(x);
61: // 'x' should contain 84 as all its entries
62: VecView(x, PETSC_VIEWER_STDOUT_WORLD);
64: VecDestroy(&x);
65: clReleaseContext(ctx);
66: clReleaseCommandQueue(queue);
67: clReleaseMemObject(mem);
68: clReleaseProgram(prg);
69: clReleaseKernel(knl);
71: 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*/