Week 3.5 - Running late

During the previous ~1.5 weeks I've been working on finishing the requirements of milestone 1 (i.e. get_image_* builtins), which means I've greatly underestimated the time required to complete it. Fortunately I've allocated some buffer time on the last weeks of GSoC to implement some optional requirements, so I'll probably have to sacrifice some of that. Now that that's out of the way, I can talk about the current state of the get_image_* builtins along with the choices I've made while implementing these features.

The general idea is essentially the same as the one I described in my previous post: the getters have to be replaced by an LLVM pass with implicit kernel parameters loads. However the code had to be restructured because several parts were in wrong places, e.g. one should avoid patching the IR in the driver if possible. To this end the pass has been moved to LLVM.

The current implementation can be broken down to 3 parts:

• builtin definitions using dummy intrinsics (libclc)
• translation of dummy intrinsics to meaningful code (LLVM)
• placement of image attribute data to specific locations before kernel launch (mesa)

The libclc definitions are simple functions which contain calls to the llvm.AMDGPU.get.image.[23]d dummy intrinsics. A new pass (R600ImageAttributeIntrinsicsReplacer) has been added to the AMDGPU backend of LLVM to replace calls to these dummy intrinsics with the newly added llvm.AMDGPU.read.image.attribute intrinsic. This intrinsic accepts two compile-time constant operands (each 4 bytes wide): an image index (i) and an attribute index (j). Upon instruction lowering, the intrinsic will be translated to a load from an implicit kernel parameter using the 4-byte offset 4 + 5 * i + j added to the starting location of the implicit parameters (the first four 4-bytes are used by grid dim and grid offset, and there are five 4-byte image attributes for each image). The attribute index is 0 for width, 1 for height, 2 for depth, 3 for channel data type and 4 for channel order. The image index is the index of the image argument among all image arguments. This decision has some consequences though. Namely it affects how convenient it is to prepare implicit arguments for the software component which is responsible for that.

This raises the question which component of the driver should prepare these implicit arguments? Currently clover serializes grid dim and grid offset taking care of byte extension and endian conversion, but ideally the pipe driver should make the choice how it implements mechanisms like grid dim info or image attributes. On the other hand image attributes like channel order and data type contain OpenCL specific constants, which the pipe driver stores in a different format. Adding conversion code to the driver would mean (1) adding state tracker specific code to the driver and (2) duplicating some functionality, so one has to maintain both the OpenCL to pipe format converter in clover and the pipe format to OpenCL converter in the driver.

So what implications do the choice above has on image indexing? The driver has no immediate knowledge of the order of read-only and write-only kernel arguments, since those are handled differently at driver level. If one chooses to prepare image attributes in the driver, another indexing scheme has to be introduced, e.g. by copying the attributes of write-only images first and the read-only ones after that.

Considering the complications arising from preparing image attributes in the driver, I decided to do it in clover instead in a similar fashion to how grid dim and grid offset are handled. Because of that, image attributes can be added to the kernel parameters simply in the order of the image arguments as is present in the kernel signature. This also simplifies the intrinsic replacer pass.

Patches have been sent to the relevant mailing list under the following subjects:

• [Mesa-dev] [PATCH] clover: Implement image attribute getters (commit on GitHub)
• [llvm] Image attribute access for the AMDGPU backend (commit on GitHub)
• [Libclc-dev] [PATCH] Add image attribute getter builtins (commit on GitHub)

UPDATE: added links to commits on GitHub.

Week 2 - Milestone 1

This week I managed to implement the OpenCL image attribute getter builtins for 2D images (and partially for 3D). These builtins allow the query of image metadata - width, height, depth (for 3D), channel data type and channel order - within an OpenCL kernel. In the remainder of this post I will describe how the implementation works and how did I got there.

Back in April when I wrote my GSoC proposal, my mentor Tom Stellard suggested that examining the ISA emitted by the Catalyst driver could help in the design process, so I did just that. I used the simple OpenCL kernel below:

__kernel void imgtest_basic(read_only image2d_t img, __global int * out)
{
*out = get_image_width(img);
}

and obtained the following ISA using AMD's CodeXL on Windows with a Mobility Radeon HD 5850 (Juniper, Evergreen) GPU:

; --------  Disassembly --------------------

00 ALU: ADDR(32) CNT(3) KCACHE0(CB1:0-15) 

      0  x: LSHR        R1.x,  KC0[2].x,  2      

         t: MOV         R0.x,  KC0[0].x      

01 MEM_RAT_CACHELESS_STORE_RAW: RAT(11)[R1].x___, R0, ARRAY_SIZE(4)  VPM 

END_OF_PROGRAM

What I think is important (or what I managed to understand) here is that the width of the image is stored in KC0[0].x which references a constant cache. The location of the data is known at compile time, no pointer dereferencing is needed. Furthermore the value of the out pointer is stored in a similar manner - in KC0[2].x, so one may assume that kernel arguments are passed this way.
Note: global addresses are 4-byte aligned on this architecture, so the least significant 2 bits should be omitted from the load addresses. This is why a right shift (LSHR) by 2 is required to transform the address stored in the argument out to a valid load input.

Clover, a software component of mesa which manages OpenCL state (a state tracker in gallium parlance), compiles OpenCL C sources to ISA using clang and LLVM. The OpenCL builtins are implemented as an LLVM bitcode library (libclc), which gets linked to the kernel in IR form. The fact that LLVM IR represents the image[23]d_t type as an opaque pointer imposes some restrictions on the implementation of the getter builtins, because this pointer is the sole input which these builtins may use to do their thing. And you can't do much with a pointer without dereferencing it.

What all this rambling lead us is that a code transformation is required to achieve something similar to what Catalyst does. So I've implemented an LLVM pass which replaces calls to the getter builtins with accesses to hidden kernel arguments containing the image attributes. The builtins are implemented in libclc as nonexistent intrinsics to mark the places where the pass needs to perform a substitution, and also to avoid problems with mangled symbol names created because of overloading. For each explicit image argument (only 2d for now) 5 new ones are created for the 5 attributes in a specific order. The LLVM pass assumes the hidden argumens are in this order when performing the substitutions. The hidden arguments will receive their values when the image argument is bound to the execution context.

A patch is on the way, I'll link it here if it gets accepted.

UPDATE:
I got some feedback for the patches. This feature needs some work, because I've added the new LLVM pass, which implements a driver-specific behavior (i.e. passes image attributes as hidden arguments), to clover, which should manage OpenCL state in a driver-independent way. So I'll have to move this to a driver-specific place. The plan is to generate intrinsics in a similar fashion to the workitem id and size getter builtins, and lower those intrinsics in the R600 LLVM backend to reads from implicit arguments.

UPDATE #2: I've got the address-shifting-by-2 thing backwards: a _right_ shift is needed before loads. See the updated note above.

Week 1 - onwards to Milestone 1

I'm halfway through the deadline of Milestone 1, which requires me to be able to create an image, pass it to a kernel and query its attributes by June 8. These are the steps I've done towards these requirements this week:

• Image support capability queries now return true. This makes it possible to call clCreateImage*.
• Added dummy get_image_width to libclc for 2d images, which now always returns 42.
• Enabled PIPE_BIND_COMPUTE_RESOURCE usage in the evergreen_is_format_supported function. Currently it returns true to any input (TODO: find out if this is correct).
• Depth value is now 1 instead of 0, because libdrm requires it (see first comment).

This is the kernel I'm testing with right now, my first goal is to be able to run this, and get back 42 in the output buffer.

__kernel void imgtest(__read_only image2d_t in, __global int * out)
{
    if (get_global_id(0) == 0) {
        *out = get_image_width(in);
    }
}

What remains to be done next week:

• Allow kernel launch to set dummy compute sampler views (currently an assert prevents that).
• Pass the image attributes to the kernel as hidden parameters, and actually use them in libclc.

Google Summer of Code project accepted

My Google Summer of Code project proposal titled OpenCL image support for the r600 gallium driver has been accepted. See the proposal below:

• Contact: Zoltan Gilian <zoltan.gilian@gmail.com>, zogi on IRC
• Title: OpenCL image support for the r600 gallium driver
• Synopsis:
  Current open source radeon drivers do not support OpenCL images. The aim of this project is to implement support for image-related OpenCL C 1.1 capabilities, both 2d and 3d, in the r600 gallium driver.
• Benefits to the Community:
  OpenCL image handling functions could utilize appropriate GPU resources (e.g. texture sampling units), therefore increasing the performance of related algorithms (e.g. image processing). Such software relying on OpenCL image capabilities cannot run using the open source drivers despite having a capable GPU in the system. Implementing OpenCL image support would solve this problem.
• Deliverables:
  • Expose image support to clover. Required.
  • Implement 2d image attribute query and basic data access OpenCL C builtins. Required.
    • querying image attributes (get_image_* builtins)
      reading image data with basic sampling parameters (using unnormalized coordinates, ‘none’ addressing mode and nearest filtering)
    • writing image data
  • Implement image sampling and access qualifiers. Optional.
  • Support 3d images. Optional.
  • Support half floating-point image access. Optional.
  • Documentation. Required.
• Milestones:
  • Be able create an image, pass it to a kernel and query its attributes by June 8.
    Functions to impement:
    • get_image_width
    • get_image_height
    • get_image_channel_data_type
    • get_image_channel_order
    • get_image_dim
  • Be able to write and read 1 dimensional image data using basic sampling parameters by June 26.
    Functions to implement:
    • read_image{f,i,ui} (with the restrictions above)
    • write_image{f,i,ui} (with the restrictions above)
  • Implement 2d image access using basic sampling parameters by July 13.
    Functions to implement:
    • read_image{f,i,ui} (with the restrictions above)
    • write_image{f,i,ui} (with the restrictions above)
  • Document the previous steps to future developers by July 20.
  • Implement and document (a subset of) the optional deliverables in the remaining time.
    • Image sampling
      • extend the functionality of read_image
    • 3d images
      • extend the functionality of the builtins above
      • new function: get_image_depth
    • Access qualifiers (read_only and write_only)
    • Half floating-point image access
      • read_imageh, write_imageh
• Architecture overview:
  • LLVM frontend (libclc): extend the libclc llvm bitcode library by supporting image related OpenCL C constructs (see: Deliverables).
    • Image attributes can be implemented as hidden kernel parameters. These attributes should be transferred to constant memory on the device.
  • LLVM backend (llvm): support intrinsics emitted by libclc.
  • r600 pipe driver (mesa): expose image support to clover.
• Related Work:
  • http://portablecl.org: OpenCL C frontend using clang.
• Biographical Information:
  • Education: MSc degree in Computer Science, Eotvos Lorand University, Budapest (2012). Currently pursuing PhD in signal processing there.
  • Work experience:
    • SRE intern at Google, Zurich (2013 Aug-Dec)
    • Intern and coop student at Morgan Stanley, Budapest (2010 summer, 2011 spring)
    • System administrator at Bolyai College, Budapest (since 2014)
  • Related technical skills: C, C++, CUDA (udacity parallel computing class) and OpenCL, Linux, compilers (BSc thesis about a compiler extension to the feldspar language, 2010, sadly in hungarian).
  • Open source work:
    • Implemented hardware page accounting into hplip cups backend for college printer accounting (hplip fork)
    • Small open source spectrum analyzer in javascript: https://github.com/zogi/spectrum-analyzer-js (demo)