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: PetscErrorCode ierr;
23: PetscInt size=5;
24: Vec x;
25: cl_program prg;
26: cl_kernel knl;
27: PETSC_UINTPTR_T clctxptr;
28: PETSC_UINTPTR_T clqueueptr;
29: PETSC_UINTPTR_T clmemptr;
30: const size_t gsize=10, lsize=2;
32: PetscInitialize(&argc,&argv,(char*)0,help);if (ierr) return ierr;
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 ierr;
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*/