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, NULL, 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*/