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!

Japan Meteorological Agency Takes Delivery of Pair of Crays

May 21, 2018

Cray has supplied two identical Cray XC50 supercomputers to the Japan Meteorological Agency (JMA) in northwestern Tokyo. Boasting more than 18 petaflops combined peak computing capacity, the new systems will extend the a Read more…

By Tiffany Trader

ASC18: Final Results Revealed & Wrapped Up

May 17, 2018

It was an exciting week at ASC18 in Nanyang, China. The student teams braved extreme heat, extremely difficult applications, and extreme competition in order to cross the cluster competition finish line. The gala awards ceremony took place on Wednesday. The auditorium was packed with student teams, various dignitaries, the media, and other interested parties. So what happened? Read more…

By Dan Olds

ASC18: Tough Applications & Tough Luck

May 17, 2018

The applications at the ASC18 Student Cluster Competition were tough. Tougher than the $3.99 steak special at your local greasy spoon restaurant. The apps are so tough that even Chuck Norris backs away from them slowly. Read more…

By Dan Olds

HPE Extreme Performance Solutions

HPC and AI Convergence is Accelerating New Levels of Intelligence

Data analytics is the most valuable tool in the digital marketplace – so much so that organizations are employing high performance computing (HPC) capabilities to rapidly collect, share, and analyze endless streams of data. Read more…

IBM Accelerated Insights

Mastering the Big Data Challenge in Cognitive Healthcare

Patrick Chain, genomics researcher at Los Alamos National Laboratory, posed a question in a recent blog: What if a nurse could swipe a patient’s saliva and run a quick genetic test to determine if the patient’s sore throat was caused by a cold virus or a bacterial infection? Read more…

Spring Meetings Underscore Quantum Computing’s Rise

May 17, 2018

The month of April 2018 saw four very important and interesting meetings to discuss the state of quantum computing technologies, their potential impacts, and the technology challenges ahead. These discussions happened in Read more…

By Alex R. Larzelere

Japan Meteorological Agency Takes Delivery of Pair of Crays

May 21, 2018

Cray has supplied two identical Cray XC50 supercomputers to the Japan Meteorological Agency (JMA) in northwestern Tokyo. Boasting more than 18 petaflops combine Read more…

By Tiffany Trader

ASC18: Final Results Revealed & Wrapped Up

May 17, 2018

It was an exciting week at ASC18 in Nanyang, China. The student teams braved extreme heat, extremely difficult applications, and extreme competition in order to cross the cluster competition finish line. The gala awards ceremony took place on Wednesday. The auditorium was packed with student teams, various dignitaries, the media, and other interested parties. So what happened? Read more…

By Dan Olds

Spring Meetings Underscore Quantum Computing’s Rise

May 17, 2018

The month of April 2018 saw four very important and interesting meetings to discuss the state of quantum computing technologies, their potential impacts, and th Read more…

By Alex R. Larzelere

Quantum Network Hub Opens in Japan

May 17, 2018

Following on the launch of its Q Commercial quantum network last December with 12 industrial and academic partners, the official Japanese hub at Keio University is now open to facilitate the exploration of quantum applications important to science and business. The news comes a week after IBM announced that North Carolina State University was the first U.S. university to join its Q Network. Read more…

By Tiffany Trader

Democratizing HPC: OSC Releases Version 1.3 of OnDemand

May 16, 2018

Making HPC resources readily available and easier to use for scientists who may have less HPC expertise is an ongoing challenge. Open OnDemand is a project by t Read more…

By John Russell

PRACE 2017 Annual Report: Exascale Aspirations; Industry Collaboration; HPC Training

May 15, 2018

The Partnership for Advanced Computing in Europe (PRACE) today released its annual report showcasing 2017 activities and providing a glimpse into thinking about Read more…

By John Russell

US Forms AI Brain Trust

May 11, 2018

Amid calls for a U.S. strategy for promoting AI development, the Trump administration is forming a senior-level panel to help coordinate government and industry research efforts. The Select Committee on Artificial Intelligence was announced Thursday (May 10) during a White House summit organized by the Office of Science and Technology Policy (OSTP). Read more…

By George Leopold

Emerging Advanced Scale Tech Trends Focus of Annual Tabor Conference

May 9, 2018

At Tabor Communications' annual Advanced Scale Forum (ASF) held this week in Austin, the focus was on enterprise adoption of HPC-class technologies and high performance data analytics (HPDA). It’s a confab that brings together end users (CIOs, IT planners, department heads) and vendors and encourages... Read more…

By the Editorial Team

MLPerf – Will New Machine Learning Benchmark Help Propel AI Forward?

May 2, 2018

Let the AI benchmarking wars begin. Today, a diverse group from academia and industry – Google, Baidu, Intel, AMD, Harvard, and Stanford among them – releas Read more…

By John Russell

How the Cloud Is Falling Short for HPC

March 15, 2018

The last couple of years have seen cloud computing gradually build some legitimacy within the HPC world, but still the HPC industry lies far behind enterprise I Read more…

By Chris Downing

Russian Nuclear Engineers Caught Cryptomining on Lab Supercomputer

February 12, 2018

Nuclear scientists working at the All-Russian Research Institute of Experimental Physics (RFNC-VNIIEF) have been arrested for using lab supercomputing resources to mine crypto-currency, according to a report in Russia’s Interfax News Agency. Read more…

By Tiffany Trader

Nvidia Responds to Google TPU Benchmarking

April 10, 2017

Nvidia highlights strengths of its newest GPU silicon in response to Google's report on the performance and energy advantages of its custom tensor processor. Read more…

By Tiffany Trader

Deep Learning at 15 PFlops Enables Training for Extreme Weather Identification at Scale

March 19, 2018

Petaflop per second deep learning training performance on the NERSC (National Energy Research Scientific Computing Center) Cori supercomputer has given climate Read more…

By Rob Farber

AI Cloud Competition Heats Up: Google’s TPUs, Amazon Building AI Chip

February 12, 2018

Competition in the white hot AI (and public cloud) market pits Google against Amazon this week, with Google offering AI hardware on its cloud platform intended Read more…

By Doug Black

US Plans $1.8 Billion Spend on DOE Exascale Supercomputing

April 11, 2018

On Monday, the United States Department of Energy announced its intention to procure up to three exascale supercomputers at a cost of up to $1.8 billion with th Read more…

By Tiffany Trader

Lenovo Unveils Warm Water Cooled ThinkSystem SD650 in Rampup to LRZ Install

February 22, 2018

This week Lenovo took the wraps off the ThinkSystem SD650 high-density server with third-generation direct water cooling technology developed in tandem with par Read more…

By Tiffany Trader

Leading Solution Providers

HPC and AI – Two Communities Same Future

January 25, 2018

According to Al Gara (Intel Fellow, Data Center Group), high performance computing and artificial intelligence will increasingly intertwine as we transition to Read more…

By Rob Farber

Researchers Measure Impact of ‘Meltdown’ and ‘Spectre’ Patches on HPC Workloads

January 17, 2018

Computer scientists from the Center for Computational Research, State University of New York (SUNY), University at Buffalo have examined the effect of Meltdown Read more…

By Tiffany Trader

Google Chases Quantum Supremacy with 72-Qubit Processor

March 7, 2018

Google pulled ahead of the pack this week in the race toward "quantum supremacy," with the introduction of a new 72-qubit quantum processor called Bristlecone. Read more…

By Tiffany Trader

HPE Wins $57 Million DoD Supercomputing Contract

February 20, 2018

Hewlett Packard Enterprise (HPE) today revealed details of its massive $57 million HPC contract with the U.S. Department of Defense (DoD). The deal calls for HP Read more…

By Tiffany Trader

CFO Steps down in Executive Shuffle at Supermicro

January 31, 2018

Supermicro yesterday announced senior management shuffling including prominent departures, the completion of an audit linked to its delayed Nasdaq filings, and Read more…

By John Russell

Deep Learning Portends ‘Sea Change’ for Oil and Gas Sector

February 1, 2018

The billowing compute and data demands that spurred the oil and gas industry to be the largest commercial users of high-performance computing are now propelling Read more…

By Tiffany Trader

Nvidia Ups Hardware Game with 16-GPU DGX-2 Server and 18-Port NVSwitch

March 27, 2018

Nvidia unveiled a raft of new products from its annual technology conference in San Jose today, and despite not offering up a new chip architecture, there were still a few surprises in store for HPC hardware aficionados. Read more…

By Tiffany Trader

Hennessy & Patterson: A New Golden Age for Computer Architecture

April 17, 2018

On Monday June 4, 2018, 2017 A.M. Turing Award Winners John L. Hennessy and David A. Patterson will deliver the Turing Lecture at the 45th International Sympo Read more…

By Staff

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