{"id":1369,"date":"2025-10-17T15:22:06","date_gmt":"2025-10-17T12:22:06","guid":{"rendered":"https:\/\/paragroup.cse.uoi.gr\/wpsite\/?page_id=1369"},"modified":"2025-12-15T16:16:22","modified_gmt":"2025-12-15T14:16:22","slug":"ompi-offloading-to-opencl-gpus","status":"publish","type":"page","link":"https:\/\/paragroup.cse.uoi.gr\/wpsite\/software\/ompi\/documentation\/ompi-offloading-to-opencl-gpus\/","title":{"rendered":"OMPi: Offloading to OpenCL GPUs"},"content":{"rendered":"

Introduction<\/h2>\n

The latest version of OMPi extends GPU support by adding the ability to offload computations to any GPU through OpenCL<\/a>. OMPi analyzes the OpenMP target directives appearing in a user application and generates kernel source files written in the OpenCL C language; these are then compiled into OpenCL kernels, and get linked\u00a0with an OpenCL library that offers OpenMP facilities during kernel execution. OpenCL support is provided by the new OMPi\u00a0opencl<\/code> module.<\/p>\n

Requirements<\/h2>\n

The requirements for the opencl<\/code>\u00a0module to be activated are the following:<\/p>\n

    \n
  • The standard OpenCL header and development packages must be installed (usually named opencl-headers, opencl-dev or ocl-icd-devel<\/code>)<\/li>\n
  • A\u00a0GPU that supports OpenCL version \u2265 1.2; you can always double-check this by running the\u00a0clinfo<\/code> utility, if it is installed on your system.<\/li>\n
  • The corresponding vendor drivers and\/or relevant libraries must be installed (we cannot provide help on how to install them as this differs from vendor to vendor;\u00a0 Intel\u00a0drivers most probably can be found in the intel-neo<\/code> repository; for AMD you need the amdgpu<\/code>\u00a0and\/or ROCm<\/code> drivers;\u00a0 NVIDIA includes the OpenCL drivers\/libraries in their CUDA SDK).<\/li>\n<\/ul>\n

    Installing OMPi with OpenCL support<\/h2>\n

    You do not need any special preparations; simply deploy OMPi on your system as usual:<\/p>\n

    meson setup build --prefix=<install-dir>\r\ncd build\/\r\nmeson compile\r\nmeson install<\/pre>\n
    If all the requirements listed above are met, the installation process will automatically detect the GPU(s) and build the opencl<\/code>\u00a0module; if something is missing, the module won’t be included. You can verify the correct installation of the opencl<\/code> module by executing:<\/p>\n
    ompiconf --devvinfo<\/pre>\n
    which lists all the identified modules\/devices, along with their numeric device IDs<\/em>. If the above command fails to show an OpenCL-capable GPU, double check that all requirements are met and ensure you have installed OMPi correctly.<\/p>\n
    <\/span>Sample output of devvinfo<\/div>
    \n
    The ompiconf --devvinfo<\/code> command (or \u00a0ompiconf --devvinfo=opencl<\/code>) should give an output similar to the following (here an integrated Intel UHD770 GPU was detected by OMPi):<\/p>\n
    MODULE [opencl]: \r\n------ \r\nOMPi module for OpenCL devices \r\nAvailable devices : 1 \r\n \r\ndevice id < 0 > { \u00a0\r\n \u00a0GPU device name \u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0: Intel(R) UHD Graphics 770 \r\n \u00a0GPU device vendor \u00a0\u00a0\u00a0\u00a0\u00a0: Intel(R) Corporation \r\n \u00a0Vendor device driver \u00a0\u00a0: 24.22.29735.20 \r\n \u00a0OpenCL ver. (platform) : 3.0 \r\n \u00a0OpenCL ver. (device) \u00a0\u00a0: 3.0 \r\n \u00a0OpenCL ver. (compiler) : 1.0, 1.1, 1.2, 3.0 \r\n \u00a0Num of compute units \u00a0\u00a0: 32 \r\n \u00a0Max workgroup size \u00a0\u00a0\u00a0\u00a0: 512 \r\n \u00a0Local memory \u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0: 64 KBytes \r\n \u00a0Global memory \u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0\u00a0: 13.9 GBytes \r\n \u00a0Coarse grain SVM \u00a0\u00a0\u00a0\u00a0\u00a0\u00a0: yes \r\n \u00a0Fine grain buffer SVM \u00a0: no \r\n \u00a0Fine grain system SVM \u00a0: no \r\n \u00a0Unified host memory \u00a0\u00a0\u00a0: yes \r\n \u00a0Supports \"double\" type : no \r\n \r\n \u00a0Intel GPU details: \r\n \u00a0\u00a0\u00a0| 1 slice(s) \r\n \u00a0\u00a0\u00a0| 2 subslice(s) per slice (max) \r\n \u00a0\u00a0\u00a0| 16 EUs per subslice (max) \r\n \u00a0\u00a0\u00a0| 7 threads per EU (max) \r\n} \r\n------\r\n<\/div><\/div><\/pre>\n
    Quick start<\/h2>\n
    To offload to a GPU through OpenCL, you need to compile your applications with a\u00a0--devs=opencl<\/code>\u00a0argument:<\/p>\n
    ompicc --devs=opencl app.c<\/pre>\n
    <\/span>Note<\/div>
    \n
    If you use the --devs=opencl<\/code>\u00a0compilation option, the numeric device IDs are those printed by\u00a0ompiconf --devinfo=opencl<\/code>\u00a0. For example, in our system we get:\u00a0<\/p>\n
    MODULE [opencl] provides device(s) : 0 1<\/code><\/p>\n
    You\u00a0may\u00a0instead compile using the\u00a0--devs=all<\/code>\u00a0option to produce kernels that may run on any GPU target, including<\/em> OpenCL.\u00a0Be careful that in this case the numeric device IDs may be different depending on what other modules\u00a0are active. In the case you use\u00a0--devs=all<\/code>, the\u00a0correct device IDs can be found by exectuting\u00a0ompiconf --devinfo<\/code>. In our system it reported:<\/p>\n
    MODULE [cuda] provides device(s) : 0<\/code>
    \n  MODULE [opencl] provides device(s) : 1 2<\/code><\/p>\n
    Thus the numeric IDs of the two OpenCL-capable GPUs are now 1 and 2.<\/p>\n<\/div><\/div>\n
    The compiler will produce the main application executable (a.out<\/code>) and a number of OpenCL kernel files, one for each\u00a0target<\/code> construct in the\u00a0OpenMP application.<\/p>\n
    <\/span>Sample OpenMP application utilizing the GPU<\/div>
    \n
    Here is a sample application that can be used to verify the\u00a0successful deployment of the opencl<\/code>\u00a0module:<\/p>\n
    #include<\/strong> <stdio.h>\r\n#include<\/strong> <omp.h>\r\n\r\nint<\/strong> main(void<\/strong>)\r\n{\r\n    \/* \r\n     * Expected result (assuming that the GPU device ID is 0):\r\n     * Running on OpenCL device \r\n     *\/\r\n    int<\/strong> x;\r\n    #pragma omp target map(tofrom:x) device(0)<\/span><\/strong>\r\n    {<\/span>\r\n        x = <\/span>omp_is_initial_device<\/span>();<\/span>\r\n    }\r\n\r\n<\/span>    if<\/span> <\/strong>(<\/span>x<\/span>)<\/span> \r\n        printf<\/span>(<\/span>\"Running on host<\/span>\\n<\/span>\"<\/span>);<\/span>    \r\n    else<\/span> <\/strong>\r\n        printf<\/span>(<\/span>\"Running on an OpenCL device\\n\");\r\n<\/span>    \r\n    return<\/strong> 0; \r\n} <\/span><\/pre>\n<\/div><\/div>\n
    <\/span>Advanced usage: Speeding up compilation<\/div>
    \n
    When an\u00a0application contains several target<\/code> regions, compilation times may become substantial. One can generate multiple compilation jobs that handle the compilation of different kernels in parallel, by using the\u00a0-j<number><\/code> device option. For example to have 4 parallel kernel compilation jobs, use ompicc<\/code> as follows:<\/p>\n
    ompicc --devs=all --devopt -j4<\/strong> app.c<\/pre>\n<\/div><\/div>\n
    Limitations<\/h2>\n
    Currently, the opencl<\/code>\u00a0module has a number of limitations:<\/p>\n
      \n
    1. Only the datatypes supported by OpenCL C can be used; some GPUs may lack\u00a0double<\/code>\u00a0support.<\/li>\n
    2. Only static loop scheduling is supported in target teams distribute parallel for<\/code> directives,<\/li>\n
    3. Stand-alone parallel<\/code>\u00a0constructs are not handled yet; the parallel<\/code>\u00a0construct must be part of a combined \u00a0target<\/code>,\u00a0teams<\/code>\u00a0and\/or\u00a0distribute<\/code>\u00a0construct.<\/li>\n
    4. OpenCL\u00a0\u2265 1.2 is required; in addition the\u00a0OpenCL C driver\/compiler must allow __global<\/code> declarations at\u00a0file scope; while this requirement\u00a0is not part of the OpenCL 1.2 standard (it is even optional in OpenCL 2.0), all major vendors support it.<\/li>\n<\/ol>\n
      Notes<\/h2>\n
      The OMPi opencl<\/code>\u00a0module has been tested with the following GPUs:<\/p>\n
        \n
      • Integrated Intel GPUs (UHD 770, Iris Xe)<\/li>\n
      • AMD Radeon GPUs (R9 285, RX 550)<\/li>\n
      • NVIDIA GPUs (GeForce GT730, Tesla P40, Ampere A2)<\/li>\n<\/ul>\n","protected":false},"excerpt":{"rendered":"
        Introduction The latest version of OMPi extends GPU support by adding the ability to offload computations to any GPU through OpenCL. OMPi analyzes the OpenMP target directives appearing in a user application and generates kernel source files written in the OpenCL C language; these are then compiled into OpenCL kernels, and get linked\u00a0with an OpenCL […]<\/p>\n","protected":false},"author":4,"featured_media":0,"parent":120,"menu_order":0,"comment_status":"closed","ping_status":"closed","template":"","meta":{"footnotes":""},"class_list":["post-1369","page","type-page","status-publish","hentry"],"_links":{"self":[{"href":"https:\/\/paragroup.cse.uoi.gr\/wpsite\/wp-json\/wp\/v2\/pages\/1369","targetHints":{"allow":["GET"]}}],"collection":[{"href":"https:\/\/paragroup.cse.uoi.gr\/wpsite\/wp-json\/wp\/v2\/pages"}],"about":[{"href":"https:\/\/paragroup.cse.uoi.gr\/wpsite\/wp-json\/wp\/v2\/types\/page"}],"author":[{"embeddable":true,"href":"https:\/\/paragroup.cse.uoi.gr\/wpsite\/wp-json\/wp\/v2\/users\/4"}],"replies":[{"embeddable":true,"href":"https:\/\/paragroup.cse.uoi.gr\/wpsite\/wp-json\/wp\/v2\/comments?post=1369"}],"version-history":[{"count":17,"href":"https:\/\/paragroup.cse.uoi.gr\/wpsite\/wp-json\/wp\/v2\/pages\/1369\/revisions"}],"predecessor-version":[{"id":1428,"href":"https:\/\/paragroup.cse.uoi.gr\/wpsite\/wp-json\/wp\/v2\/pages\/1369\/revisions\/1428"}],"up":[{"embeddable":true,"href":"https:\/\/paragroup.cse.uoi.gr\/wpsite\/wp-json\/wp\/v2\/pages\/120"}],"wp:attachment":[{"href":"https:\/\/paragroup.cse.uoi.gr\/wpsite\/wp-json\/wp\/v2\/media?parent=1369"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}