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 industry updates delivered to you every week!

Q&A with Nvidia’s Chief of DGX Systems on the DGX-GB200 Rack-scale System

March 27, 2024

Pictures of Nvidia's new flagship mega-server, the DGX GB200, on the GTC show floor got favorable reactions on social media for the sheer amount of computing power it brings to artificial intelligence.  Nvidia's DGX Read more…

Call for Participation in Workshop on Potential NSF CISE Quantum Initiative

March 26, 2024

Editor’s Note: Next month there will be a workshop to discuss what a quantum initiative led by NSF’s Computer, Information Science and Engineering (CISE) directorate could entail. The details are posted below in a Ca Read more…

Waseda U. Researchers Reports New Quantum Algorithm for Speeding Optimization

March 25, 2024

Optimization problems cover a wide range of applications and are often cited as good candidates for quantum computing. However, the execution time for constrained combinatorial optimization applications on quantum device Read more…

NVLink: Faster Interconnects and Switches to Help Relieve Data Bottlenecks

March 25, 2024

Nvidia’s new Blackwell architecture may have stolen the show this week at the GPU Technology Conference in San Jose, California. But an emerging bottleneck at the network layer threatens to make bigger and brawnier pro Read more…

Who is David Blackwell?

March 22, 2024

During GTC24, co-founder and president of NVIDIA Jensen Huang unveiled the Blackwell GPU. This GPU itself is heavily optimized for AI work, boasting 192GB of HBM3E memory as well as the the ability to train 1 trillion pa Read more…

Nvidia Appoints Andy Grant as EMEA Director of Supercomputing, Higher Education, and AI

March 22, 2024

Nvidia recently appointed Andy Grant as Director, Supercomputing, Higher Education, and AI for Europe, the Middle East, and Africa (EMEA). With over 25 years of high-performance computing (HPC) experience, Grant brings a Read more…

Q&A with Nvidia’s Chief of DGX Systems on the DGX-GB200 Rack-scale System

March 27, 2024

Pictures of Nvidia's new flagship mega-server, the DGX GB200, on the GTC show floor got favorable reactions on social media for the sheer amount of computing po Read more…

NVLink: Faster Interconnects and Switches to Help Relieve Data Bottlenecks

March 25, 2024

Nvidia’s new Blackwell architecture may have stolen the show this week at the GPU Technology Conference in San Jose, California. But an emerging bottleneck at Read more…

Who is David Blackwell?

March 22, 2024

During GTC24, co-founder and president of NVIDIA Jensen Huang unveiled the Blackwell GPU. This GPU itself is heavily optimized for AI work, boasting 192GB of HB Read more…

Nvidia Looks to Accelerate GenAI Adoption with NIM

March 19, 2024

Today at the GPU Technology Conference, Nvidia launched a new offering aimed at helping customers quickly deploy their generative AI applications in a secure, s Read more…

The Generative AI Future Is Now, Nvidia’s Huang Says

March 19, 2024

We are in the early days of a transformative shift in how business gets done thanks to the advent of generative AI, according to Nvidia CEO and cofounder Jensen Read more…

Nvidia’s New Blackwell GPU Can Train AI Models with Trillions of Parameters

March 18, 2024

Nvidia's latest and fastest GPU, codenamed Blackwell, is here and will underpin the company's AI plans this year. The chip offers performance improvements from Read more…

Nvidia Showcases Quantum Cloud, Expanding Quantum Portfolio at GTC24

March 18, 2024

Nvidia’s barrage of quantum news at GTC24 this week includes new products, signature collaborations, and a new Nvidia Quantum Cloud for quantum developers. Wh Read more…

Houston We Have a Solution: Addressing the HPC and Tech Talent Gap

March 15, 2024

Generations of Houstonian teachers, counselors, and parents have either worked in the aerospace industry or know people who do - the prospect of entering the fi Read more…

Alibaba Shuts Down its Quantum Computing Effort

November 30, 2023

In case you missed it, China’s e-commerce giant Alibaba has shut down its quantum computing research effort. It’s not entirely clear what drove the change. Read more…

Nvidia H100: Are 550,000 GPUs Enough for This Year?

August 17, 2023

The GPU Squeeze continues to place a premium on Nvidia H100 GPUs. In a recent Financial Times article, Nvidia reports that it expects to ship 550,000 of its lat Read more…

Shutterstock 1285747942

AMD’s Horsepower-packed MI300X GPU Beats Nvidia’s Upcoming H200

December 7, 2023

AMD and Nvidia are locked in an AI performance battle – much like the gaming GPU performance clash the companies have waged for decades. AMD has claimed it Read more…

DoD Takes a Long View of Quantum Computing

December 19, 2023

Given the large sums tied to expensive weapon systems – think $100-million-plus per F-35 fighter – it’s easy to forget the U.S. Department of Defense is a Read more…

Synopsys Eats Ansys: Does HPC Get Indigestion?

February 8, 2024

Recently, it was announced that Synopsys is buying HPC tool developer Ansys. Started in Pittsburgh, Pa., in 1970 as Swanson Analysis Systems, Inc. (SASI) by John Swanson (and eventually renamed), Ansys serves the CAE (Computer Aided Engineering)/multiphysics engineering simulation market. Read more…

Choosing the Right GPU for LLM Inference and Training

December 11, 2023

Accelerating the training and inference processes of deep learning models is crucial for unleashing their true potential and NVIDIA GPUs have emerged as a game- Read more…

Intel’s Server and PC Chip Development Will Blur After 2025

January 15, 2024

Intel's dealing with much more than chip rivals breathing down its neck; it is simultaneously integrating a bevy of new technologies such as chiplets, artificia Read more…

Baidu Exits Quantum, Closely Following Alibaba’s Earlier Move

January 5, 2024

Reuters reported this week that Baidu, China’s giant e-commerce and services provider, is exiting the quantum computing development arena. Reuters reported � Read more…

Leading Solution Providers

Contributors

Comparing NVIDIA A100 and NVIDIA L40S: Which GPU is Ideal for AI and Graphics-Intensive Workloads?

October 30, 2023

With long lead times for the NVIDIA H100 and A100 GPUs, many organizations are looking at the new NVIDIA L40S GPU, which it’s a new GPU optimized for AI and g Read more…

Shutterstock 1179408610

Google Addresses the Mysteries of Its Hypercomputer 

December 28, 2023

When Google launched its Hypercomputer earlier this month (December 2023), the first reaction was, "Say what?" It turns out that the Hypercomputer is Google's t Read more…

AMD MI3000A

How AMD May Get Across the CUDA Moat

October 5, 2023

When discussing GenAI, the term "GPU" almost always enters the conversation and the topic often moves toward performance and access. Interestingly, the word "GPU" is assumed to mean "Nvidia" products. (As an aside, the popular Nvidia hardware used in GenAI are not technically... Read more…

Shutterstock 1606064203

Meta’s Zuckerberg Puts Its AI Future in the Hands of 600,000 GPUs

January 25, 2024

In under two minutes, Meta's CEO, Mark Zuckerberg, laid out the company's AI plans, which included a plan to build an artificial intelligence system with the eq Read more…

Google Introduces ‘Hypercomputer’ to Its AI Infrastructure

December 11, 2023

Google ran out of monikers to describe its new AI system released on December 7. Supercomputer perhaps wasn't an apt description, so it settled on Hypercomputer Read more…

China Is All In on a RISC-V Future

January 8, 2024

The state of RISC-V in China was discussed in a recent report released by the Jamestown Foundation, a Washington, D.C.-based think tank. The report, entitled "E Read more…

Intel Won’t Have a Xeon Max Chip with New Emerald Rapids CPU

December 14, 2023

As expected, Intel officially announced its 5th generation Xeon server chips codenamed Emerald Rapids at an event in New York City, where the focus was really o Read more…

IBM Quantum Summit: Two New QPUs, Upgraded Qiskit, 10-year Roadmap and More

December 4, 2023

IBM kicks off its annual Quantum Summit today and will announce a broad range of advances including its much-anticipated 1121-qubit Condor QPU, a smaller 133-qu Read more…

  • arrow
  • Click Here for More Headlines
  • arrow
HPCwire