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