Official OpenCL SPIR 1.2 feedback thread

OpenCL SPIR 1.2 Provisional Specification
The OpenCL Working Group released the OpenCL SPIR 1.2 provisional specification for public review. ‘SPIR’ stands for Standard Portable Intermediate Representation and is a portable non-source representation for OpenCL 1.2 device programs. It enables application developers to avoid shipping kernel source and to manage the proliferation of devices and drivers from multiple vendors. OpenCL SPIR will enable consumption of code from third party compiler front-ends for alternative languages, such as C++, and is based on LLVM 3.2. Khronos has contributed open source patches for Clang 3.2 to enable SPIR code generation.

OpenCL BOF at SIGGRAPH, Anaheim, CA July 24th 2013
There is an OpenCL BOF “Birds of a Feather” Meeting on Wednesday July 24th at 4-5PM at the Hilton Anaheim, California Ballroom A & B, where attendees are invited to meet OpenCL implementers and developers and learn more about the new OpenCL 2.0 specification.

[ul]
[li]OpenCL BOF[/li][li]OpenCL homepage[/li][/ul]

Maybe it’s my reading, but I am missing the information/specification how to access the Thread geometry/Id in SPIR 1.2 provisional. Are LLVM intrinstics going to be declared for things like

  • get_global_id()
  • get_local_id()
  • get_num_groups()
  • get_group_id()

Frank

In the SPIR 1.2 spec ( section 1.2 Name mangling and Apendix A SPIR name mangling ), you should auto-rename the static/inlined functions and variables to protect better the kernel’s IP. For instance, a function like this:

static vec4 testAABBHitRay ( vec4 pmin, vec4 pmax, vec4 o, vec4 dir )
{
__local float rayHits[2048];
}

should be re-named in a compiler’s pre-pass as:

static vec4 func1 ( vec4 arg1, vec4 arg2, vec4 arg3, vec4 arg4 )
{
__local float localVar1[2048];
}

so it’s less visible in the final SPIR / PTX code.

Also, this will help with the text’s chars longitude to be processed by the compiler.

Have the SPIR specificaton version match the OpenCL specification version it is part of.
Thus in this case change the SPIR version to 2.0. IF you need a version to denote api compatibility do this with a seperate something that indicates what lowest api version it is compatible with.

There doesn’t seem mention of C’s fixed with integer types (found in inttypes.h) in the specification.
It would be practical to have OpenCL, OpenCL SPIR cover that too.
(Types are described as fixed width integers.)
An unambiguous mapping would ensure interoperability, compatibility with code that uses these types.

Hi. I have written some feedback for SPIR on my blog. I have to admit that I do not have extensive experience working on a compiler, but I have addressed concerns that I have regarding future-proofing something like SPIR, along with the role that I think SPIR shoul play in OpenCL (which isn’t entirely in agreement with the SPIR spec). Please take a look, and comment on my blog (or here).

https://blog.ajguillon.com/2013/10/08/opencl-2-0-spir-feedback-and-vision/