[cfe-dev] OpenCL support - using metadata

[Benyei, Guy]

Hi Anton,
You can see the current metadata we've below.
Basically, it holds a list of the kernels, as in your patch, but for each kernel it also holds some additional data in the following order: required work group size (or (0,0,0) if not defined), work group size hint (or (0,0,0) if not defined), vec type hint (as a string), the kernel's signature (as a string too), and the name of a named metadata node, that holds a list of the kernel's locals' names.
We'd also like to know how your metadata works - I'm sure we can find a common solution that will be agreed on everyone.
BTW, local variables in our solution are treated much like "static" C variables; they are promoted to the program scope with prefixes (as you might see below). According to my understanding, function scope constant variables are prohibited, the OpenCL 1.1 spec says at 6.5.3: "Variables allocated in the __constant address space can only be defined as program scope variables and are required to be initialized."


kernel __attribute__(( vec_type_hint(char2) )) void foo(global char4* in, global int16* out)
{
  local float4 ff4;
  size_t id = get_global_id(0);
  out[id] = (int16)(as_int(in[id]));
}

kernel __attribute__(( work_group_size_hint(16,1,1) )) void foo2(global char4* in, global int16* out)
{
  local float4 ff4;
  size_t id = get_global_id(0);
  out[id] = (int16)(as_int(in[id]));
}



!opencl.kernels = !{!0, !2}
!opencl_foo_locals_anchor = !{!4}
!opencl_foo2_locals_anchor = !{!5}

!0 = metadata !{void (<4 x i8> addrspace(1)*, <16 x i32> addrspace(1)*)* @foo, metadata !1, metadata !1, metadata !"char2", metadata !"char4 __attribute__((address_space(1))) *, int16 __attribute__((address_space(1))) *", metadata !"opencl_foo_locals_anchor"}
!1 = metadata !{i32 0, i32 0, i32 0}
!2 = metadata !{void (<4 x i8> addrspace(1)*, <16 x i32> addrspace(1)*)* @foo2, metadata !1, metadata !3, metadata !"", metadata !"char4 __attribute__((address_space(1))) *, int16 __attribute__((address_space(1))) *", metadata !"opencl_foo2_locals_anchor"}
!3 = metadata !{i32 16, i32 1, i32 1}
!4 = metadata !{metadata !"opencl_foo_local_ff4"}
!5 = metadata !{metadata !"opencl_foo2_local_ff4"}


Thanks
      Guy Benyei
      Intel
      SSG - MGP OpenCL Development Center



-----Original Message-----
From: Anton Lokhmotov [mailto:Anton.Lokhmotov at arm.com] 
Sent: Thursday, February 24, 2011 17:59
To: Benyei, Guy
Cc: cfe-dev at cs.uiuc.edu
Subject: Re: OpenCL support - using metadata

Hi Guy,

I would like to gather more information about how different implementations
represent OpenCL C programs in LLVM-IR in order to collate the best practice
and make it into a standard.
 
> We at Intel are also using metadata for our implementation. This
> approach is very useful for a few reasons:
> First, when compiling OpenCL C code to llvm, some important information
> about the kernel arguments is lost. i.e. typedefs are lost, structures
> are sometimes disassembled. We're adding a string containing the
> interesting part original function signature.
Hmm, if we want to standardise this, we need to define which parts of the
signature are considered "interesting".

> Second, we use a global metadata node, which enumerates all the kernel
> functions. This can be used to quickly enumerate the kernels in a
> program, instead of passing the whole module.
We are doing exactly the same: 
http://lists.cs.uiuc.edu/pipermail/cfe-dev/2010-December/012741.html

> Third, for each kernel we hold the kernel attributes in the same
> metadata: vector type hint, required workgroup size and workgroup size
> hint.
Do you use one metadata node per kernel?  (It seems we are doing it
differently.)  Do these nodes have the same format (three fields per
kernel)?  What do you use for defaults?

> Last, we also hold in the metadata a list of the local variables
> defined in that kernel.
Do you mean kernel-scope variables declared in the local address space?  How
about kernel-scope variables declared in the constant address space?   We
are bringing both classes to the program scope, prefixing them with the
kernel's name.  Do you have any views on that?

Best, Anton.



---------------------------------------------------------------------
Intel Israel (74) Limited

This e-mail and any attachments may contain confidential material for
the sole use of the intended recipient(s). Any review or distribution
by others is strictly prohibited. If you are not the intended
recipient, please contact the sender and delete all copies.