OpenCL: To GPGPU and Beyond

By John West

December 11, 2008

This week at SIGGRAPH Asia in Singapore, The Khronos Group announced that version 1.0 of the OpenCL specification has been ratified. In the short term OpenCL is expected to encourage the development of applications that can take advantage of those GPGPUs you can find in just about every machine these days (though not in the way you might expect). But in the long term the implications for HPC may be much more far-reaching than “just” the GPU.

OpenCL originated inside the hallowed halls of Steve Jobs’ Apple – it will be integrated into Snow Leopard, the next release of the company’s operating system, as part of the suite of technologies being introduced to streamline Mac OS X and improve the performance of applications. In June of 2008 Apple turned a draft standard over to The Khronos Group, “a member-funded industry consortium focused on the creation of open standard, royalty-free APIs to enable the authoring and accelerated playback of dynamic media on a wide variety of platforms and devices.” You may not have heard of Khronos, but they are the team that is responsible for the care and feeding of OpenGL, and several other less well known standards.

This week, only six months after it started work, Khronos has announced that the OpenCL consortium has reached consensus on version 1.0 of the specification. That seems pretty impressive, especially when you consider the long list of IT heavyweights who were involved: Intel, AMD, Ericsson, IBM, NVIDIA, Nokia, Texas Instruments, Apple, Motorola (and more; see the full list). OpenCL is based on C99, a dialect of C that was adopted as an ANSI standard in May 2000.

But what’s it do? You’ve no doubt heard that OpenCL will accelerate the development of applications that take advantage of the extra processing power available in the nearly ubiquitous GPU. Both AMD and NVIDIA have announced support for the standard alongside their own Brook+ and CUDA efforts. For example, any card that supports CUDA will support OpenCL. And Intel is a member of the consortium, so the conventional wisdom is that support is expected for Larrabee as well, although Intel has been pretty vague so far while it actively develops its own language for throughput computing, Ct.

OpenCL is an important step in spurring the development of application for GPUs, because developers would like to accelerate their code on all available GPUs rather than locking them — and their customers — into one vendor’s solution. However, OpenCL is a low-level standard when compared with other alternatives. For example, programmers will have to do their own memory management (unlike with CUDA), and overall OpenCL supports fewer productivity-enhancing abstractions. In the Khronos Group’s own slide deck OpenCL is characterized as “approachable, but primarily targeted at expert developers.” This means that OpenCL’s benefits may come to application developers indirectly, through the tools community. Companies like Codeplay and RapidMind are members of the consortium and will likely incorporate the technology into their development platforms.

What about the 800-pound gorilla OS? Microsoft is not part of the consortium, and is in fact developing a competing technology called DirectX 11 Compute. NVIDIA plans to provide support for this effort as well (in an Engadget article NVIDIA says “it will go where the customers are.”) But because the GPUs vendors themselves are planning to support OpenCL directly, Windows developers will be able to get at the specification, or use tools that take advantage of it, whether or not Microsoft officially endorses OpenCL.

OpenCL may prove pretty significant for traditional HPC as well as the commodity user crowd. OpenCL may well provide ISPs an avenue for delivering high levels of performance on traditional scientific computing applications for users of servers and desktops by encouraging more applications to take advantage of the GPU. OpenCL is much more ambitious than just the GPU, however, as it aims to provide a “programming environment for software developers to write efficient, portable code for high-performance compute servers, desktop computer systems and handheld devices using a diverse mix of multicore CPUs, GPUs, Cell-type architectures and other parallel processors such as DSPs.” In fact, The Khronos Group’s web site and briefing materials make frequent mention of HPC, and it is clear that the consortium members have big things in mind for the new standard as it relates to HPC.

SGI’s Bob Pette, the vice president in charge of visualization at the company, sees a real need in the marketplace for the kinds of issues addressed by OpenCL. “The biggest issue we hear about from customers using accelerator and multi-paradigm computing is overcoming the programming challenges,” he says. “The other major concern is protection of their software investment as hardware technology change rapidly. OpenCL is a positive step forward in addressing both of these critical problems and should fuel growth and advancement in utilization of multicore and manycore technologies.”

Multicore CPUs are an important target for HPC, especially for those interested in using low-end HPC or developing technologies to promote the widespread adoption of HPC. OpenCL, or OpenCL-based tools, will enable developers to parallelize an application for a system with multicore processors, GPUs, or both, amplifying the possibilities for high performance computation on commodity hardware. Another nice benefit of OpenCL’s support for both multicore and GPU solutions is that ISVs don’t have to decide for themselves whether, as some companies are betting, GPUs will return to irrelevance for most computation as the number of cores available on processing chips increases.

OpenCL supports both data- and task-parallel programming models. Task-parallel programming doesn’t really fit well on GPUs but is applicable to the Cell processor or Larrabee. This portability will provide developers a single abstraction with the ability to write applications that can run on everything from a run-of-the-mill PC from WalMart to IBM’s RoadRunner (although some target-specific code may be needed for performance reasons, which we won’t know until the specification gets some performance testing in the real world).

OpenCL does impose some restrictions on the programming model. For example, recursion is a no-no, pointers to functions aren’t allowed, and pointers themselves are only allowed within a kernel, not as arguments. Bit fields aren’t supported, and neither are variable length arrays or structures (see slide 38 of the Khronos deck for a more complete discussion of restrictions). Interestingly the language provides a shared memory model with “relaxed consistency,” and implementations map whatever physical memories are available into the private/local/global address spaces of the specification.

So, what does an OpenCL program look like? I’m glad you asked. The OpenCL slide deck referenced above provides two example implementations, a parallel n-body implementation and a simple vector addition. Since the vector addition is conceptually more straightforward, let’s look at it here (this example starts on slide 43, and the slides have pointers to the specification for more details on specific concepts if you really want to dive into this).

The problem is to compute c = a + b, where a, b, and c are vectors of length N. The kernel that is called by the main program is straightforward:

    __kernel void vec_add (__global const float *a,
    __global const float *b,
    __global float *c)
    {
        int gid = get_global_id(0);
        c[gid] = a[gid] + b[gid];
    }

The main program isn’t particularly complicated, but it is surprisingly long, and it has some fairly odd-looking constructs that set up the context, allocate the memory, and then invoke the kernel we just looked at.

The first step is to query the platform and create (in this case) a GPU compute context and the command queue:

    // create the OpenCL context on a GPU device
    cl_context context = clCreateContextFromType(0, // (must be 0)
                    CL_DEVICE_TYPE_GPU,
                    NULL, // error callback
                    NULL, // user data
                    NULL); // error code

    // get the list of GPU devices associated with context
    size_t cb;
    clGetContextInfo(context, CL_CONTEXT_DEVICES, 0, NULL, &cb);
    cl_device_id *devices = malloc(cb);
    clGetContextInfo(context, CL_CONTEXT_DEVICES, cb, devices, NULL);
   
    // create a command-queue
    cl_cmd_queue cmd_queue = clCreateCommandQueue(context,
                    devices[0],
                    0, // default options
                    NULL); // error code

Next, we create memory objects to hold arrays a, b, and c:

    cl_mem memobjs[3];

    // allocate input buffer memory objects
    memobjs[0] = clCreateBuffer(context,
                    CL_MEM_READ_ONLY | // flags
                    CL_MEM_COPY_HOST_PTR,
                    sizeof(cl_float)*n, // size
                    srcA, // host pointer
                    NULL); // error code

    memobjs[1] = clCreateBuffer(context,
                    CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
                    sizeof(cl_float)*n, srcB, NULL);

    // allocate input buffer memory object
    memobjs[2] = clCreateBuffer(context, CL_MEM_WRITE_ONLY,
                    sizeof(cl_float)*n, NULL, NULL);

After this, we create and build the program and create (we aren’t calling yet) the kernel:

    // create the program
    cl_program program = clCreateProgramWithSource(
                    context,
                    1, // string count
                    &program_source, // program strings
                    NULL, // string lengths
                    NULL); // error code

    // build the program
    cl_int err = clBuildProgram(program,
                    0, // num devices in device list
                    NULL, // device list
                    NULL, // options
                    NULL, // notifier callback function ptr
                    NULL); // error code

    // create the kernel
    cl_kernel kernel = clCreateKernel(program, “vec_add”, NULL);

I was on board until this last section, especially the creating and building of the program within the program, but for those more familiar with the use of accelerators this code may not look that odd.

Finally we set the kernel arguments to point to a, b, and c:

    // set vector argument 0
    err = clSetKernelArg(kernel,
                    0, // argument index
                    (void *)&memobjs[0], // argument data
                    sizeof(cl_mem)); // argument data size

    // set  vector argument 1
    err |= clSetKernelArg(kernel, 1, (void *)&memobjs[1], sizeof(cl_mem));

    // set vector argument 2
    err |= clSetKernelArg(kernel, 2, (void *)&memobjs[2], sizeof(cl_mem));

Invoke the kernel, and read the results:

    size_t global_work_size[1] = n; // set work-item dimensions

    // execute kernel
    err = clEnqueueNDRangeKernel(cmd_queue, kernel,
                    1, // Work dimensions
                    NULL, // must be NULL (work offset)
                    global_work_size,
                    NULL, // automatic local work size
                    0, // no events to wait on
                    NULL, // event list
                    NULL); // event for this kernel

    // read output array
    err = clEnqueueReadBuffer( context, memobjs[2],
                    CL_TRUE, // blocking
                    0, // offset
                    n*sizeof(cl_float), // size
                    dst, // pointer
                    0, NULL, NULL); // events

All in all, this looks like a lot of code to add a couple vectors in parallel, and this relates to the points made earlier about OpenCL being “approachable” to the regular programmer, but really made for the “expert programmer in you” (with apologies to Frosted Mini-Wheats). Going through the n-body example in the slides (which, if you are still reading, I recommend) shows more of the power of the language relative to the computational tasks being performed.

My feeling is that OpenCL is a significant development for the broader application development community, spanning as it does everything from cell phones to servers. It also appears to have a lot of potential for supporting the democratization of HPC by encouraging ISVs to make the investment to improve the performance of their applications without having to make potentially costly bets on which, and whose, processing technology will dominate over the next 5 to 7 years. As SGI’s Pette points out, there is still a lot of work to be done to realize the potential of OpenCL, “There will likely be a performance penalty — seeing as the major vendors may choose to layer OpenCL on top of their native drivers. And of course, mere ratification does not imply implementation.”

Subscribe to HPCwire's Weekly Update!

Be the most informed person in the room! Stay ahead of the tech trends with industy updates delivered to you every week!

HPC Startup Advances Auto-Parallelization’s Promise

January 23, 2017

The shift from single core to multicore hardware has made finding parallelism in codes more important than ever, but that hasn’t made the task of parallel programming any easier. Read more…

By Tiffany Trader

Answered Prayers for High Frequency Traders? Latency Cut to 20 Nanoseconds

January 23, 2017

“You can buy your way out of bandwidth problems. But latency is divine.”

This sentiment, from Intel Technical Computing Group CTO Mark Seager, seems as old as the Bible, a truth universally acknowledged. Read more…

By Doug Black

CMU’s Latest “Card Shark” – Libratus – is Beating the Poker Pros (Again)

January 20, 2017

It’s starting to look like Carnegie Mellon University has a gambling problem – can’t stay away from the poker table. Read more…

By John Russell

IDG to Be Bought by Chinese Investors; IDC to Spin Out HPC Group

January 19, 2017

US-based publishing and investment firm International Data Group, Inc. (IDG) will be acquired by a pair of Chinese investors, China Oceanwide Holdings Group Co., Ltd. Read more…

By Tiffany Trader

HPE Extreme Performance Solutions

Enhancing Patient Care with Next-Generation Sequencing

In the ever-evolving world of life sciences, speed, accuracy, and savings are more important than ever. Today’s scientists and healthcare professionals are leveraging high-performance computing (HPC) solutions to solve the world’s greatest health problems and accelerate the diagnoses and treatment of a variety of medical conditions. Read more…

Weekly Twitter Roundup (Jan. 19, 2017)

January 19, 2017

Here at HPCwire, we aim to keep the HPC community apprised of the most relevant and interesting news items that get tweeted throughout the week. Read more…

By Thomas Ayres

France’s CEA and Japan’s RIKEN to Partner on ARM and Exascale

January 19, 2017

France’s CEA and Japan’s RIKEN institute announced a multi-faceted five-year collaboration to advance HPC generally and prepare for exascale computing. Among the particulars are efforts to: build out the ARM ecosystem; work on code development and code sharing on the existing and future platforms; share expertise in specific application areas (material and seismic sciences for example); improve techniques for using numerical simulation with big data; and expand HPC workforce training. It seems to be a very full agenda. Read more…

By Nishi Katsuya and John Russell

ARM Waving: Attention, Deployments, and Development

January 18, 2017

It’s been a heady two weeks for the ARM HPC advocacy camp. At this week’s Mont-Blanc Project meeting held at the Barcelona Supercomputer Center, Cray announced plans to build an ARM-based supercomputer in the U.K. while Mont-Blanc selected Cavium’s ThunderX2 ARM chip for its third phase of development. Last week, France’s CEA and Japan’s Riken announced a deep collaboration aimed largely at fostering the ARM ecosystem. This activity follows a busy 2016 when SoftBank acquired ARM, OpenHPC announced ARM support, ARM released its SVE spec, Fujistu chose ARM for the post K machine, and ARM acquired HPC tool provider Allinea in December. Read more…

By John Russell

Women Coders from Russia, Italy, and Poland Top Study

January 17, 2017

According to a study posted on HackerRank today the best women coders as judged by performance on HackerRank challenges come from Russia, Italy, and Poland. Read more…

By John Russell

HPC Startup Advances Auto-Parallelization’s Promise

January 23, 2017

The shift from single core to multicore hardware has made finding parallelism in codes more important than ever, but that hasn’t made the task of parallel programming any easier. Read more…

By Tiffany Trader

Answered Prayers for High Frequency Traders? Latency Cut to 20 Nanoseconds

January 23, 2017

“You can buy your way out of bandwidth problems. But latency is divine.”

This sentiment, from Intel Technical Computing Group CTO Mark Seager, seems as old as the Bible, a truth universally acknowledged. Read more…

By Doug Black

IDG to Be Bought by Chinese Investors; IDC to Spin Out HPC Group

January 19, 2017

US-based publishing and investment firm International Data Group, Inc. (IDG) will be acquired by a pair of Chinese investors, China Oceanwide Holdings Group Co., Ltd. Read more…

By Tiffany Trader

France’s CEA and Japan’s RIKEN to Partner on ARM and Exascale

January 19, 2017

France’s CEA and Japan’s RIKEN institute announced a multi-faceted five-year collaboration to advance HPC generally and prepare for exascale computing. Among the particulars are efforts to: build out the ARM ecosystem; work on code development and code sharing on the existing and future platforms; share expertise in specific application areas (material and seismic sciences for example); improve techniques for using numerical simulation with big data; and expand HPC workforce training. It seems to be a very full agenda. Read more…

By Nishi Katsuya and John Russell

ARM Waving: Attention, Deployments, and Development

January 18, 2017

It’s been a heady two weeks for the ARM HPC advocacy camp. At this week’s Mont-Blanc Project meeting held at the Barcelona Supercomputer Center, Cray announced plans to build an ARM-based supercomputer in the U.K. while Mont-Blanc selected Cavium’s ThunderX2 ARM chip for its third phase of development. Last week, France’s CEA and Japan’s Riken announced a deep collaboration aimed largely at fostering the ARM ecosystem. This activity follows a busy 2016 when SoftBank acquired ARM, OpenHPC announced ARM support, ARM released its SVE spec, Fujistu chose ARM for the post K machine, and ARM acquired HPC tool provider Allinea in December. Read more…

By John Russell

Spurred by Global Ambitions, Inspur in Joint HPC Deal with DDN

January 17, 2017

Inspur, the fast-growth cloud computing and server vendor from China that has several systems on the current Top500 list, and DDN, a leader in high-end storage, have announced a joint sales and marketing agreement to produce solutions based on DDN storage platforms integrated with servers, networking, software and services from Inspur. Read more…

By Doug Black

For IBM/OpenPOWER: Success in 2017 = (Volume) Sales

January 11, 2017

To a large degree IBM and the OpenPOWER Foundation have done what they said they would – assembling a substantial and growing ecosystem and bringing Power-based products to market, all in about three years. Read more…

By John Russell

UberCloud Cites Progress in HPC Cloud Computing

January 10, 2017

200 HPC cloud experiments, 80 case studies, and a ton of hands-on experience gained, that’s the harvest of four years of UberCloud HPC Experiments. Read more…

By Wolfgang Gentzsch and Burak Yenier

AWS Beats Azure to K80 General Availability

September 30, 2016

Amazon Web Services has seeded its cloud with Nvidia Tesla K80 GPUs to meet the growing demand for accelerated computing across an increasingly-diverse range of workloads. The P2 instance family is a welcome addition for compute- and data-focused users who were growing frustrated with the performance limitations of Amazon's G2 instances, which are backed by three-year-old Nvidia GRID K520 graphics cards. Read more…

By Tiffany Trader

For IBM/OpenPOWER: Success in 2017 = (Volume) Sales

January 11, 2017

To a large degree IBM and the OpenPOWER Foundation have done what they said they would – assembling a substantial and growing ecosystem and bringing Power-based products to market, all in about three years. Read more…

By John Russell

US, China Vie for Supercomputing Supremacy

November 14, 2016

The 48th edition of the TOP500 list is fresh off the presses and while there is no new number one system, as previously teased by China, there are a number of notable entrants from the US and around the world and significant trends to report on. Read more…

By Tiffany Trader

Vectors: How the Old Became New Again in Supercomputing

September 26, 2016

Vector instructions, once a powerful performance innovation of supercomputing in the 1970s and 1980s became an obsolete technology in the 1990s. But like the mythical phoenix bird, vector instructions have arisen from the ashes. Here is the history of a technology that went from new to old then back to new. Read more…

By Lynd Stringer

Container App ‘Singularity’ Eases Scientific Computing

October 20, 2016

HPC container platform Singularity is just six months out from its 1.0 release but already is making inroads across the HPC research landscape. It's in use at Lawrence Berkeley National Laboratory (LBNL), where Singularity founder Gregory Kurtzer has worked in the High Performance Computing Services (HPCS) group for 16 years. Read more…

By Tiffany Trader

Dell EMC Engineers Strategy to Democratize HPC

September 29, 2016

The freshly minted Dell EMC division of Dell Technologies is on a mission to take HPC mainstream with a strategy that hinges on engineered solutions, beginning with a focus on three industry verticals: manufacturing, research and life sciences. "Unlike traditional HPC where everybody bought parts, assembled parts and ran the workloads and did iterative engineering, we want folks to focus on time to innovation and let us worry about the infrastructure," said Jim Ganthier, senior vice president, validated solutions organization at Dell EMC Converged Platforms Solution Division. Read more…

By Tiffany Trader

D-Wave SC16 Update: What’s Bo Ewald Saying These Days

November 18, 2016

Tucked in a back section of the SC16 exhibit hall, quantum computing pioneer D-Wave has been talking up its new 2000-qubit processor announced in September. Forget for a moment the criticism sometimes aimed at D-Wave. This small Canadian company has sold several machines including, for example, ones to Lockheed and NASA, and has worked with Google on mapping machine learning problems to quantum computing. In July Los Alamos National Laboratory took possession of a 1000-quibit D-Wave 2X system that LANL ordered a year ago around the time of SC15. Read more…

By John Russell

Enlisting Deep Learning in the War on Cancer

December 7, 2016

Sometime in Q2 2017 the first ‘results’ of the Joint Design of Advanced Computing Solutions for Cancer (JDACS4C) will become publicly available according to Rick Stevens. He leads one of three JDACS4C pilot projects pressing deep learning (DL) into service in the War on Cancer. Read more…

By John Russell

Leading Solution Providers

Lighting up Aurora: Behind the Scenes at the Creation of the DOE’s Upcoming 200 Petaflops Supercomputer

December 1, 2016

In April 2015, U.S. Department of Energy Undersecretary Franklin Orr announced that Intel would be the prime contractor for Aurora: Read more…

By Jan Rowell

CPU Benchmarking: Haswell Versus POWER8

June 2, 2015

With OpenPOWER activity ramping up and IBM’s prominent role in the upcoming DOE machines Summit and Sierra, it’s a good time to look at how the IBM POWER CPU stacks up against the x86 Xeon Haswell CPU from Intel. Read more…

By Tiffany Trader

Nvidia Sees Bright Future for AI Supercomputing

November 23, 2016

Graphics chipmaker Nvidia made a strong showing at SC16 in Salt Lake City last week. Read more…

By Tiffany Trader

Beyond von Neumann, Neuromorphic Computing Steadily Advances

March 21, 2016

Neuromorphic computing – brain inspired computing – has long been a tantalizing goal. The human brain does with around 20 watts what supercomputers do with megawatts. And power consumption isn’t the only difference. Fundamentally, brains ‘think differently’ than the von Neumann architecture-based computers. While neuromorphic computing progress has been intriguing, it has still not proven very practical. Read more…

By John Russell

BioTeam’s Berman Charts 2017 HPC Trends in Life Sciences

January 4, 2017

Twenty years ago high performance computing was nearly absent from life sciences. Today it’s used throughout life sciences and biomedical research. Genomics and the data deluge from modern lab instruments are the main drivers, but so is the longer-term desire to perform predictive simulation in support of Precision Medicine (PM). There’s even a specialized life sciences supercomputer, ‘Anton’ from D.E. Shaw Research, and the Pittsburgh Supercomputing Center is standing up its second Anton 2 and actively soliciting project proposals. There’s a lot going on. Read more…

By John Russell

The Exascale Computing Project Awards $39.8M to 22 Projects

September 7, 2016

The Department of Energy’s Exascale Computing Project (ECP) hit an important milestone today with the announcement of its first round of funding, moving the nation closer to its goal of reaching capable exascale computing by 2023. Read more…

By Tiffany Trader

Dell Knights Landing Machine Sets New STAC Records

November 2, 2016

The Securities Technology Analysis Center, commonly known as STAC, has released a new report characterizing the performance of the Knight Landing-based Dell PowerEdge C6320p server on the STAC-A2 benchmarking suite, widely used by the financial services industry to test and evaluate computing platforms. The Dell machine has set new records for both the baseline Greeks benchmark and the large Greeks benchmark. Read more…

By Tiffany Trader

What Knights Landing Is Not

June 18, 2016

As we get ready to launch the newest member of the Intel Xeon Phi family, code named Knights Landing, it is natural that there be some questions and potentially some confusion. Read more…

By James Reinders, Intel

  • arrow
  • Click Here for More Headlines
  • arrow
Share This