[clang][WIP] add option to keep types of ptr args for non-kernel functions in metadata

[iliya-diyachkov]

Summary

It's a draft for the patch we discussed in the thread of opaque pointer support in the SPIR-V translator. The final decision on whether we will use this approach has not yet been made, so the patch is uploaded for convenience and concertizing of the discussion.

The types are already in metadata for kernel functions by default. Alexey Bader shared the code example with non-kernel functions which lost the info. The patch fixes this case. To enable the feature you need to pass -cl-extra-ptr-info to clang.

TODO:

• emit only ptr type info metadata if EmitOpenCLExtraPtrInfo is enabled,
• maybe add the same info to function declarations,
• maybe add the same options to other languages (SYCL, HLSL ...) or make one common option for all languages.

[Anastasia]

The patch seems very straight forward. I wonder though if we should always enable this for SPIR-V w/o adding any flag until the untyped pointers are added to SPIR-V so it will become a target setting controlled instead of a flag?

Also is there anything that we don't need from the metadata that are currently emitted for the use case i.e. for example argument names? This could reduce the number of metadata emitted...

I am not sure if this is something that other languages or backends might also need in which case we might need to make this more universal... Ideally it would be nice if it was not in metadata as it is somewhat optional... but potentially we don't have many other alternatives?

[beanz]

There are numerous changes floating around that boil down to "we need pointer types". Most are SPIR-V & DXIL related. I'm a little concerned that we are generally building a bunch of different approaches to tracking pointer types and not considering a more general solution that could meet all of the needs.

[nikic]

Could you please add some context directly to the patch description that explains why this is necessary? Being unfamiliar with SPRIV, I don't really get what is significant about the code example.

[beanz]

@nikic the most important thing you need to know about SPIR-V is that it is a virtual ISA based on LLVM IR. The ISA itself encodes types for pointers just like LLVM IR would.

[nikic]

In D127579#3585516, @beanz wrote:

@nikic the most important thing you need to know about SPIR-V is that it is a virtual ISA based on LLVM IR. The ISA itself encodes types for pointers just like LLVM IR would.

Okay ... I guess my next question would be how the SPIR-V situation differs from the DXIL situation. For DXIL you have already implemented code to "guess" pointer types insofar as they are relevant -- does SPIR-V need something more than that?

Clang is only permitted to encode pointer type information if that pointer type has some kind of direct semantic meaning -- say, if the pointer element type affects the call ABI in some way. If it does not have direct semantic meaning, then deriving the pointer type based on usage (as DXIL does) and using that for emission would be right approach.

[bader]

In D127579#3585516, @beanz wrote:

@nikic the most important thing you need to know about SPIR-V is that it is a virtual ISA based on LLVM IR. The ISA itself encodes types for pointers just like LLVM IR would.

And in addition to that ISA defines types, which are not natively supported by LLVM IR e.g. image. To represent those types clang in OpenCL language mode emits a pointer to an opaque structure with special name like opencl.<ISA_type_name> (e.g. opencl.image2d_t). All ISA types, which are defined that way look the same with type-less pointers.
https://github.com/llvm/llvm-project/blob/main/clang/include/clang/Basic/OpenCLImageTypes.def

[Anastasia]

In D127579#3585537, @nikic wrote:

In D127579#3585516, @beanz wrote:

@nikic the most important thing you need to know about SPIR-V is that it is a virtual ISA based on LLVM IR. The ISA itself encodes types for pointers just like LLVM IR would.

Okay ... I guess my next question would be how the SPIR-V situation differs from the DXIL situation. For DXIL you have already implemented code to "guess" pointer types insofar as they are relevant -- does SPIR-V need something more than that?

Clang is only permitted to encode pointer type information if that pointer type has some kind of direct semantic meaning -- say, if the pointer element type affects the call ABI in some way. If it does not have direct semantic meaning, then deriving the pointer type based on usage (as DXIL does) and using that for emission would be right approach.

I guess the problem might be that we can't always derive the type, for example if we have a translation unit:

extern void foo(int * ptr);
kernel void k(int * ptr) {
  foo(ptr);
}

then if in another translation module you have something like:

void foo(int * ptr){
  *ptr = 1;
}

For this program foo will translate into SPIR-V with different function prototype because in the first unit the type can't be deduced or deduced to some default type but in the second case it is deduced to the exact type. So the functions in two modules will have two different prototypes when mapped into SPIR-V and therefore linking won't be resolved?.. while in OpenCL sources it is just the same function foo.

[nikic]

@Anastasia Thanks, that does sound like a legitimate reason to include the information. I want to double check though, does linking the modules actually fail if the functions have signatures that differ only by pointer types? At least for normal LLVM IR this would work fine, and would just result in the insertion of a bitcast during linking (and then typically the bitcast would get shifted from the called function to the call arguments later).

[jcranmer-intel]

In D127579#3585553, @bader wrote:

And in addition to that ISA defines types, which are not natively supported by LLVM IR e.g. image. To represent those types clang in OpenCL language mode emits a pointer to an opaque structure with special name like opencl.<ISA_type_name> (e.g. opencl.image2d_t). All ISA types, which are defined that way look the same with type-less pointers.
https://github.com/llvm/llvm-project/blob/main/clang/include/clang/Basic/OpenCLImageTypes.def

One of the important facts about these types is that the SPIR-V specification doesn't let you actually cast between these types and other types such as integers. (Which is an issue because I've just come across a testcase where, in opaque pointer mode, an OpenCL event type is being used in a ptrtoint to store to a variable as an i64, which I can't legally translate to SPIR-V.) Of course, this probably means that these types need to have their representation in LLVM changed entirely, but I haven't yet done the legwork to experiment in that mode.

I haven't yet tested whether or not this patch is sufficient to support all of the use cases I need for type scavenging for SPIR-V, but one of the things I do see is that there's no support for return type information, which may be a bit of an issue.

[iliya-diyachkov]

In D127579#3585411, @Anastasia wrote:

Also is there anything that we don't need from the metadata that are currently emitted for the use case i.e. for example argument names? This could reduce the number of metadata emitted...

Sure, I think we don't need all info. As I mentioned in TODO I'm going to leave only type info and remove other metadata emission. For this purpose I'll separate a new function (that emits the type info) from GenOpenCLArgMetadata.

In D127579#3586472, @jcranmer-intel wrote:

I haven't yet tested whether or not this patch is sufficient to support all of the use cases I need for type scavenging for SPIR-V, but one of the things I do see is that there's no support for return type information, which may be a bit of an issue.

I think the return type information can be added in the next version of the patch, however it should be attached to function declarations (not definitions as it's done now), right? Do you think, declarations also require information about argument types?

[Anastasia]

In D127579#3586092, @nikic wrote:

@Anastasia Thanks, that does sound like a legitimate reason to include the information. I want to double check though, does linking the modules actually fail if the functions have signatures that differ only by pointer types? At least for normal LLVM IR this would work fine, and would just result in the insertion of a bitcast during linking (and then typically the bitcast would get shifted from the called function to the call arguments later).

@nikic If I use spirv-link with two modules that have mismatching pointee type in a function parameter I get an error:

error: 0: Type mismatch on symbol "foo" between imported variable/function %6 and exported variable/function %17.

The way I understand a bitcast instruction in SPIR-V (OpBitcast in https://www.khronos.org/registry/SPIR-V/specs/unified1/SPIRV.html#_conversion_instructions) is that it can only apply to pointer types which are distinct from function types. Note that I believe that function pointers are illegal, at least we disallow them in OpenCL.

[bader]

The way I understand a bitcast instruction in SPIR-V (OpBitcast in https://www.khronos.org/registry/SPIR-V/specs/unified1/SPIRV.html#_conversion_instructions) is that it can only apply to pointer types which are distinct from function types. Note that I believe that function pointers are illegal, at least we disallow them in OpenCL.

FYI: we are experimenting with function pointers on Intel HW programmed via SPIR-V. Extension draft - https://github.com/intel/llvm/blob/sycl/sycl/doc/design/spirv-extensions/SPV_INTEL_function_pointers.asciidoc.

[Anastasia]

In D127579#3588651, @bader wrote:

The way I understand a bitcast instruction in SPIR-V (OpBitcast in https://www.khronos.org/registry/SPIR-V/specs/unified1/SPIRV.html#_conversion_instructions) is that it can only apply to pointer types which are distinct from function types. Note that I believe that function pointers are illegal, at least we disallow them in OpenCL.

FYI: we are experimenting with function pointers on Intel HW programmed via SPIR-V. Extension draft - https://github.com/intel/llvm/blob/sycl/sycl/doc/design/spirv-extensions/SPV_INTEL_function_pointers.asciidoc.

Ok, I see. That confirms I am guessing that in the standard SPIR-V binary we won't be able to apply LLVM IR's trick.

[jcranmer-intel]

In D127579#3588607, @iliya-diyachkov wrote:

I think the return type information can be added in the next version of the patch, however it should be attached to function declarations (not definitions as it's done now), right? Do you think, declarations also require information about argument types?

Function declarations are more useful to get information than function definitions. A lot of the failures I'm currently seeing in my side of things appear to be due to not getting pointer types right when calling OpenCL/SYCL builtins causing interesting things to happen later on.

There's another approach I've considered taking, which is rather than using a metadata-based approach, emitting elementtype for any pointer-valued arguments [it doesn't solve return values, though]. This does require more changes to LLVM to let such attributes work on non-intrinsic functions, but it may end up working smoother than trying to reparse C types as LLVM types.

[aeubanks]

If I'm understanding correctly, previously the LLVM pointee type was used to represent a frontend OpenCL type. An alternative to marking everything with elementtype or adding metadata, can something be done in the frontend to mangle the function name with OpenCL pointee types, then the backend translates that to the proper SPIRV types?

[Anastasia]

In D127579#3590632, @aeubanks wrote:

If I'm understanding correctly, previously the LLVM pointee type was used to represent a frontend OpenCL type. An alternative to marking everything with elementtype or adding metadata, can something be done in the frontend to mangle the function name with OpenCL pointee types, then the backend translates that to the proper SPIRV types?

We were trying to get away from relying on mangled names in the backend as it is harder to target for non-clang and non-OpenCL based languages.

[Anastasia]

In D127579#3585537, @nikic wrote:

In D127579#3585516, @beanz wrote:

Clang is only permitted to encode pointer type information if that pointer type has some kind of direct semantic meaning -- say, if the pointer element type affects the call ABI in some way. If it does not have direct semantic meaning, then deriving the pointer type based on usage (as DXIL does) and using that for emission would be right approach.

@nikic what is the way for clang to encode the pointer type? Is it just generated regularly as before or using something like elementtype hint?

[Anastasia]

Btw do we have any issue tracking this failure, I am having a feeling that the translation of special types might be related to https://github.com/llvm/llvm-project/issues/55121? I am not aware at which point libclc is expected to be linked for SPIR-V but if linking happens at SPIR-V binary level it is likely going to face the other problem I described here.

Either way we should track this problem and we probably need a solution ideally within a month or so, as we can't release clang-15 with broken SPIR-V support... potentially one solution is to enable pointer types just for SPIR-V but this is probably the least option we should consider.

[nikic]

In D127579#3588626, @Anastasia wrote:

In D127579#3586092, @nikic wrote:

@Anastasia Thanks, that does sound like a legitimate reason to include the information. I want to double check though, does linking the modules actually fail if the functions have signatures that differ only by pointer types? At least for normal LLVM IR this would work fine, and would just result in the insertion of a bitcast during linking (and then typically the bitcast would get shifted from the called function to the call arguments later).

@nikic If I use spirv-link with two modules that have mismatching pointee type in a function parameter I get an error:

error: 0: Type mismatch on symbol "foo" between imported variable/function %6 and exported variable/function %17.

The way I understand a bitcast instruction in SPIR-V (OpBitcast in https://www.khronos.org/registry/SPIR-V/specs/unified1/SPIRV.html#_conversion_instructions) is that it can only apply to pointer types which are distinct from function types. Note that I believe that function pointers are illegal, at least we disallow them in OpenCL.

Okay ... can we maybe turn this around then? Always emit function parameters as i8* and bitcast them as needed, even if it is possible to guess a better type based on the definition? (Let's ignore the image type case here, which seems to have different requirements from the rest.)

[Anastasia]

In D127579#3595461, @nikic wrote:

In D127579#3588626, @Anastasia wrote:

In D127579#3586092, @nikic wrote:

@Anastasia Thanks, that does sound like a legitimate reason to include the information. I want to double check though, does linking the modules actually fail if the functions have signatures that differ only by pointer types? At least for normal LLVM IR this would work fine, and would just result in the insertion of a bitcast during linking (and then typically the bitcast would get shifted from the called function to the call arguments later).

@nikic If I use spirv-link with two modules that have mismatching pointee type in a function parameter I get an error:

error: 0: Type mismatch on symbol "foo" between imported variable/function %6 and exported variable/function %17.

The way I understand a bitcast instruction in SPIR-V (OpBitcast in https://www.khronos.org/registry/SPIR-V/specs/unified1/SPIRV.html#_conversion_instructions) is that it can only apply to pointer types which are distinct from function types. Note that I believe that function pointers are illegal, at least we disallow them in OpenCL.

Okay ... can we maybe turn this around then? Always emit function parameters as i8* and bitcast them as needed, even if it is possible to guess a better type based on the definition? (Let's ignore the image type case here, which seems to have different requirements from the rest.)

So where would bitcasts be emitted to reconstruct the function prototypes correctly?

[nikic]

In D127579#3598306, @Anastasia wrote:

In D127579#3595461, @nikic wrote:

In D127579#3588626, @Anastasia wrote:

In D127579#3586092, @nikic wrote:

@Anastasia Thanks, that does sound like a legitimate reason to include the information. I want to double check though, does linking the modules actually fail if the functions have signatures that differ only by pointer types? At least for normal LLVM IR this would work fine, and would just result in the insertion of a bitcast during linking (and then typically the bitcast would get shifted from the called function to the call arguments later).

@nikic If I use spirv-link with two modules that have mismatching pointee type in a function parameter I get an error:

error: 0: Type mismatch on symbol "foo" between imported variable/function %6 and exported variable/function %17.

The way I understand a bitcast instruction in SPIR-V (OpBitcast in https://www.khronos.org/registry/SPIR-V/specs/unified1/SPIRV.html#_conversion_instructions) is that it can only apply to pointer types which are distinct from function types. Note that I believe that function pointers are illegal, at least we disallow them in OpenCL.

Okay ... can we maybe turn this around then? Always emit function parameters as i8* and bitcast them as needed, even if it is possible to guess a better type based on the definition? (Let's ignore the image type case here, which seems to have different requirements from the rest.)

So where would bitcasts be emitted to reconstruct the function prototypes correctly?

The bitcasts would be in the definition only, when the pointer arguments actually get used in a typed manner. The prototype would always include i8* arguments only.

[Anastasia]

In D127579#3598308, @nikic wrote:

In D127579#3598306, @Anastasia wrote:

In D127579#3595461, @nikic wrote:

In D127579#3588626, @Anastasia wrote:

In D127579#3586092, @nikic wrote:

@Anastasia Thanks, that does sound like a legitimate reason to include the information. I want to double check though, does linking the modules actually fail if the functions have signatures that differ only by pointer types? At least for normal LLVM IR this would work fine, and would just result in the insertion of a bitcast during linking (and then typically the bitcast would get shifted from the called function to the call arguments later).

@nikic If I use spirv-link with two modules that have mismatching pointee type in a function parameter I get an error:

error: 0: Type mismatch on symbol "foo" between imported variable/function %6 and exported variable/function %17.

The way I understand a bitcast instruction in SPIR-V (OpBitcast in https://www.khronos.org/registry/SPIR-V/specs/unified1/SPIRV.html#_conversion_instructions) is that it can only apply to pointer types which are distinct from function types. Note that I believe that function pointers are illegal, at least we disallow them in OpenCL.

Okay ... can we maybe turn this around then? Always emit function parameters as i8* and bitcast them as needed, even if it is possible to guess a better type based on the definition? (Let's ignore the image type case here, which seems to have different requirements from the rest.)

So where would bitcasts be emitted to reconstruct the function prototypes correctly?

The bitcasts would be in the definition only, when the pointer arguments actually get used in a typed manner. The prototype would always include i8* arguments only.

Ok, so does it mean that all pointer parameters in function definitions compiled for SPIR-V targets in clang would just be cast to the typed pointer? Then we only have untyped pointers in the declarations of functions? Perhaps I am missing something but would it not just be easier to keep the pointer types completely as it used to be, since it is going to be present nearly everywhere apart from declarations?

clang/include/clang/Basic/CodeGenOptions.def

 CODEGENOPT(EmitDeclMetadata  , 1, 0) ///< Emit special metadata indicating what
                                      ///< Decl* various IR entities came from.
                                      ///< Only useful when running CodeGen as a
                                      ///< subroutine.
 CODEGENOPT(EmitVersionIdentMetadata , 1, 1) ///< Emit compiler version metadata.
 CODEGENOPT(EmitGcovArcs      , 1, 0) ///< Emit coverage data files, aka. GCDA.
 CODEGENOPT(EmitGcovNotes     , 1, 0) ///< Emit coverage "notes" files, aka GCNO.
 CODEGENOPT(EmitOpenCLArgMetadata , 1, 0) ///< Emit OpenCL kernel arg metadata.
+CODEGENOPT(EmitOpenCLExtraPtrInfo , 1, 0) ///< Emit OpenCL extra pointer info.
 CODEGENOPT(EmulatedTLS       , 1, 0) ///< Set by default or -f[no-]emulated-tls.
 CODEGENOPT(ExplicitEmulatedTLS , 1, 0) ///< Set if -f[no-]emulated-tls is used.
 /// Embed Bitcode mode (off/all/bitcode/marker).
 ENUM_CODEGENOPT(EmbedBitcode, EmbedBitcodeKind, 2, Embed_Off)
 /// Inline asm dialect, -masm=(att|intel)
 ENUM_CODEGENOPT(InlineAsmDialect, InlineAsmDialectKind, 1, IAD_ATT)
 CODEGENOPT(ForbidGuardVariables , 1, 0) ///< Issue errors if C++ guard variables
                                         ///< are required.

clang/include/clang/Driver/Options.td

def cl_uniform_work_group_size : Flag<["-"], "cl-uniform-work-group-size">, Group<opencl_Group>, Flags<[CC1Option]>,
   HelpText<"OpenCL only. Defines that the global work-size be a multiple of the work-group size specified to clEnqueueNDRangeKernel">,
   MarshallingInfoFlag<CodeGenOpts<"UniformWGSize">>;
 def cl_no_stdinc : Flag<["-"], "cl-no-stdinc">, Group<opencl_Group>,
   HelpText<"OpenCL only. Disables all standard includes containing non-native compiler types and functions.">;
 def cl_ext_EQ : CommaJoined<["-"], "cl-ext=">, Group<opencl_Group>, Flags<[CC1Option]>,
   HelpText<"OpenCL only. Enable or disable OpenCL extensions/optional features. The argument is a comma-separated "
            "sequence of one or more extension names, each prefixed by '+' or '-'.">,
   MarshallingInfoStringVector<TargetOpts<"OpenCLExtensionsAsWritten">>;
+def cl_extra_ptr_info : Flag<["-"], "cl-extra-ptr-info">, Group<opencl_Group>, Flags<[CC1Option]>,
+  HelpText<"OpenCL only. Generate type hints for pointers.">,
+  MarshallingInfoFlag<CodeGenOpts<"EmitOpenCLExtraPtrInfo">>;
 
 def client__name : JoinedOrSeparate<["-"], "client_name">;
 def combine : Flag<["-", "--"], "combine">, Flags<[NoXarchOption, Unsupported]>;
 def compatibility__version : JoinedOrSeparate<["-"], "compatibility_version">;
 def config : Separate<["--"], "config">, Flags<[NoXarchOption]>,
   HelpText<"Specifies configuration file">;
 def config_system_dir_EQ : Joined<["--"], "config-system-dir=">, Flags<[NoXarchOption, HelpHidden]>,
   HelpText<"System directory for configuration files">;

clang/lib/CodeGen/CodeGenFunction.cpp

CodeGenFunction::DecodeAddrUsedInPrologue(llvm::Value *F,
   // Load the original pointer through the global.
   return Builder.CreateLoad(Address(GOTAddr, Int8PtrTy, getPointerAlign()),
                             "decoded_addr");
 }
 
 void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
                                                llvm::Function *Fn)
 {
-  if (!FD->hasAttr<OpenCLKernelAttr>())
+  if (!FD->hasAttr<OpenCLKernelAttr>()) {
+    if (CGM.getCodeGenOpts().EmitOpenCLExtraPtrInfo)
+      CGM.GenOpenCLArgMetadata(Fn, FD, this);
     return;
+  }
 
   llvm::LLVMContext &Context = getLLVMContext();
 
   CGM.GenOpenCLArgMetadata(Fn, FD, this);
 
   if (const VecTypeHintAttr *A = FD->getAttr<VecTypeHintAttr>()) {
     QualType HintQTy = A->getTypeHint();
     const ExtVectorType *HintEltQTy = HintQTy->getAs<ExtVectorType>();

clang/lib/Driver/ToolChains/Clang.cpp

static void RenderOpenCLOptions(const ArgList &Args, ArgStringList &CmdArgs,
   // cl-denorms-are-zero is not forwarded. It is translated into a generic flag
   // for denormal flushing handling based on the target.
   const unsigned ForwardedArguments[] = {
       options::OPT_cl_opt_disable,
       options::OPT_cl_strict_aliasing,
       options::OPT_cl_single_precision_constant,
       options::OPT_cl_finite_math_only,
       options::OPT_cl_kernel_arg_info,
+      options::OPT_cl_extra_ptr_info,
       options::OPT_cl_unsafe_math_optimizations,
       options::OPT_cl_fast_relaxed_math,
       options::OPT_cl_mad_enable,
       options::OPT_cl_no_signed_zeros,
       options::OPT_cl_fp32_correctly_rounded_divide_sqrt,
       options::OPT_cl_uniform_work_group_size
   };