Actual source code: ex56.c
petsc-3.14.1 2020-11-03
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" ;
21: int main(int argc,char **argv)
22: {
23: PetscErrorCode ierr;
24: PetscInt size=5;
25: Vec x;
26: cl_program prg;
27: cl_kernel knl;
28: PETSC_UINTPTR_T clctxptr;
29: PETSC_UINTPTR_T clqueueptr;
30: PETSC_UINTPTR_T clmemptr;
31: const size_t gsize=10, lsize=2;
33: PetscInitialize(&argc,&argv,(char*)0,help);if (ierr) return ierr;
35: VecCreate(PETSC_COMM_WORLD,&x);
36: VecSetSizes(x,size,PETSC_DECIDE);
37: VecSetType(x, VECVIENNACL);
38: VecSet(x, 42.0);
41: VecViennaCLGetCLContext(x, &clctxptr);
42: VecViennaCLGetCLQueue(x, &clqueueptr);
43: VecViennaCLGetCLMem(x, &clmemptr);
45: const cl_context ctx = ((const cl_context)clctxptr);
46: const cl_command_queue queue = ((const cl_command_queue)clqueueptr);
47: const cl_mem mem = ((const cl_mem)clmemptr);
50: prg = clCreateProgramWithSource(ctx, 1, (const char **) & kernelSrc, NULL, NULL);
51: clBuildProgram(prg, 0, NULL, NULL, NULL, NULL);
52: knl = clCreateKernel(prg, "doublify", NULL);
54: clSetKernelArg(knl, 0, sizeof(cl_mem), &mem);
55: clSetKernelArg(knl, 1, sizeof(PetscInt), &size);
57: // Launch the kernel. (gsize > size: masked execution of some work items)
58: clEnqueueNDRangeKernel(queue, knl, 1, NULL, &gsize, &lsize, 0, NULL, NULL);
59: clFinish(queue);
61: // let petsc know that device data is altered
62: VecViennaCLRestoreCLMem(x);
64: // 'x' should contain 84 as all its entries
65: VecView(x,PETSC_VIEWER_STDOUT_WORLD);
67: VecDestroy(&x);
68: clReleaseContext(ctx);
69: clReleaseCommandQueue(queue);
70: clReleaseMemObject(mem);
71: clReleaseProgram(prg);
72: clReleaseKernel(knl);
74: PetscFinalize();
75: return ierr;
76: }
78: /*TEST
80: build:
81: requires: viennacl
83: test:
84: nsize: 1
85: suffix: 1
86: args: -viennacl_backend opencl -viennacl_opencl_device_type gpu
88: TEST*/