handling of kernel struct parameters

[a2flo]

Note: cross-posted from KhronosGroup/SPIRV-LLVM#151

given the following OpenCL C code:

typedef struct {
    int val;
} test_struct;

kernel void struct_test(global int* buf, test_struct param) {
    buf[get_global_id(0)] = param.val;
}

kernel void int_test(global int* buf, int param) {
    buf[get_global_id(0)] = param;
}

resulting in the following IR (shortened for brevity):

%struct.test_struct = type { i32 }

define spir_kernel void @struct_test(i32 addrspace(1)* %buf, %struct.test_struct* %param) nounwind {
  %1 = getelementptr inbounds %struct.test_struct* %param, i64 0, i32 0
  %2 = load i32* %1, align 4, !tbaa !12
  %3 = tail call spir_func i64 @_Z13get_global_idj(i32 0) nounwind readnone
  %4 = getelementptr inbounds i32 addrspace(1)* %buf, i64 %3
  store i32 %2, i32 addrspace(1)* %4, align 4, !tbaa !12
  ret void
}

define spir_kernel void @int_test(i32 addrspace(1)* %buf, i32 %param) nounwind {
  %1 = tail call spir_func i64 @_Z13get_global_idj(i32 0) nounwind readnone
  %2 = getelementptr inbounds i32 addrspace(1)* %buf, i64 %1
  store i32 %param, i32 addrspace(1)* %2, align 4, !tbaa !12
  ret void
}

resulting in the following SPIR-V (shortened for brevity):

               OpEntryPoint Kernel %12 "struct_test"
               OpEntryPoint Kernel %25 "int_test"
               OpDecorate %5 LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import
          %2 = OpTypeInt 64 0
          %7 = OpTypeInt 32 0
         %16 = OpConstant %2 0
         %17 = OpConstant %7 0
          %3 = OpTypeVector %2 3
          %4 = OpTypePointer UniformConstant %3
          %6 = OpTypeVoid
          %8 = OpTypePointer CrossWorkgroup %7
          %9 = OpTypeStruct %7
         %10 = OpTypePointer Function %9
         %11 = OpTypeFunction %6 %8 %10
         %18 = OpTypePointer Function %7
         %24 = OpTypeFunction %6 %8 %7
          %5 = OpVariable %4 UniformConstant
         %12 = OpFunction %6 None %11
         %13 = OpFunctionParameter %8
         %14 = OpFunctionParameter %10
         %15 = OpLabel
         %19 = OpInBoundsPtrAccessChain %18 %14 %16 %17
         %20 = OpLoad %7 %19 Aligned 4
         %21 = OpLoad %3 %5 Aligned 0
         %22 = OpCompositeExtract %2 %21 0
         %23 = OpInBoundsPtrAccessChain %8 %13 %22
               OpStore %23 %20 Aligned 4
               OpReturn
               OpFunctionEnd
         %25 = OpFunction %6 None %24
         %26 = OpFunctionParameter %8
         %27 = OpFunctionParameter %7
         %28 = OpLabel
         %29 = OpLoad %3 %5 Aligned 0
         %30 = OpCompositeExtract %2 %29 0
         %31 = OpInBoundsPtrAccessChain %8 %26 %30
               OpStore %31 %27 Aligned 4
               OpReturn
               OpFunctionEnd

Is the way kernel struct parameters are handled really the correct/intended behavior?
Considering that scalar types are directly used in OpFunctionParameter/OpTypeFunction, shouldn't structs be handled the same way instead of going through an "OpTypePointer Function" indirection? Even more, doesn't this indirection say that only a pointer argument will be set/used (4 or 8 bytes), not so much a struct object (which could be any size)?
I know that the issue here is that LLVM/SPIR can only handle struct parameters as pointers, but something like that isn't specified for SPIR-V.

How to solve this?

Option 1 (preferable):
Keep it the way it is right now, but explicitly specify that kernel pointer parameters to Function/private memory actually perform some kind of allocation of the element/pointee type on the device side, and are set as this element/pointee type on the host side (not as the pointer type). Note that private address space pointer kernel arguments are otherwise invalid.

Option 2 (impossible?):
Directly use OpTypeStruct in OpFunctionParameter/OpTypeFunction. This will however require IR/SPIR-V translator changes, since OpTypeStruct is no longer a pointer type (making all GEPs/Op*AccessChain instructions using it invalid). This might be impossible to do though, since there is no way of getting a pointer to this struct then in SPIR-V (afaik).

edit:
Option 3:
Require a OpVariable in OpFunctionParameter/OpTypeFunction for struct types. This way it should be clear what is actually happening + it is still a pointer.

[bader]

Is the way kernel struct parameters are handled really the correct/intended behavior?
Considering that scalar types are directly used in OpFunctionParameter/OpTypeFunction, shouldn't structs be handled the same way instead of going through an "OpTypePointer Function" indirection? Even more, doesn't this indirection say that only a pointer argument will be set/used (4 or 8 bytes), not so much a struct object (which could be any size)?
I know that the issue here is that LLVM/SPIR can only handle struct parameters as pointers, but something like that isn't specified for SPIR-V.

I think this is part of the ABI and not SPIR-V specification. We probably should document this behavior in OpenCL C -> SPIR-V compiler documentation, since this is specific to OpenCL C front-end compiler (aka clang), which implements C calling convention during compilation to LLVM IR.

Will this work for you?

[a2flo]

It's both, I would say. Even though this code was generated from OpenCL C, I wouldn't say it's necessarily an OpenCL problem (though OpenCL vendor runtimes/compilers should handle this properly of course). The SPIR-V code could have been generated from anything, but it would still be the same problem, i.e. struct/aggregate parameters can only be specified as pointers.

So, unless this can be solved differently, the SPIR-V spec should at the very least state that entry point function parameters of type "OpTypePointer Function" to "OpTypeStruct" are to be treated/allocated as full objects by the implementation, not as "will set sizeof(pointer) bytes". Similar to what the SPIR spec says (in 3.5 SPIR ABI):

Each function argument and return type is classified as follows:

• Any aggregate type is passed as a pointer. Memory allocation (if needed) is the responsibility of the caller function.

The "caller" being the SPIR-V/OpenCL implementation here.

[dneto0]

Interesting. I agree with @bader that this is really about what ABI is accepted in code consumed by an OpenCL implementation or anything that runs Kernel entry points. It's great that SPIR 1.2 described enough of the ABI, and perhaps we need that level of detail in the OpenCL 2.1 and 2.2 environment specs, e.g. https://www.khronos.org/registry/cl/specs/opencl-2.2-environment.pdf

Note that OpenCL C 6.9, first bullet point says: "Arguments to kernel functions declared in a program that are pointers must be declared with the __global, __constant or __local qualifier. "
I couldn't see a prohibition against using a struct argument to a kernel. But if you assume an ABI in which structs are passed as pointers, then pointer-to-private (as in your test case) would be disallowed.

I'll take this to the SPIR working group for discussion.

[sheredom]

@a2flo Can I query that the IR generated was exactly what came out of the tool? (I'm unfamiliar with the tool myself!)

The reason I ask is that I would expect a byval to be on the %struct.test_struct* %param member, and then the converter should be able to pick up the byval and know that this should become an OpTypeStruct parameter rather than a pointer to the struct.

[a2flo]

In this case, yes, it does add the byval attribute. But unless it is actually specified somewhere that this attribute is required, I'd consider all attributes to be optional and this should work regardless. And I'm not sure if LLVM will always add the byval attribute. Though looking through TargetInfo.cpp, it actually seems like this will be the case when using the default ABI (it won't emit byval when targeting x86/x64 (for sycl I guess?)). Also note that this behavior will change with LLVM 3.7+ when using C++ aggregates with non-trivial destructors/copy-constructors (https://github.com/llvm-mirror/clang/blob/release_37/lib/CodeGen/TargetInfo.cpp#L416).

But even if there is a byval attribute on the struct parameter, this will only add a "OpDecorate %14 FuncParamAttr ByVal", not transform the struct pointer to an actual struct (which would lead to the problems mentioned in option #2).

If a byval attribute for struct parameters would be mandatory in the spec, that could be a solution as well.

[sheredom]

@a2flo thanks for your response! I'll bring this up within the SPIR-V working group for discussion.

[raunintc]

@a2flo - Thanks for the feedback.

The working group agrees with @bader and @dneto0. We're currently working on describing the ABI in the OpenCL environment spec.

As @dneto0 stated, OpenCL-C does not permit private pointers as kernel arguments and our intent was to keep maintain that restriction in SPIR-V binaries used for OpenCL.

For structs, this has the added benefit that the clSetKernelArg() validation is trivial for implementations. clSetKernelArg needs to validate the size of the argument being set. Passing a struct by private pointer in SPIR-V, but by value in the host API, opens further complications that the compiler must be able to determine the size of the struct, which could be non-trivial.

The goal of SPIR-V is to directly (when possible) represent the source program. The example source program has a struct in the parameter list. I would therefore expect OpTypeStruct to show up in the parameter list as well.

[a2flo]

The working group agrees with @bader and @dneto0. We're currently working on describing the ABI in the OpenCL environment spec.

Okay, good.

As @dneto0 stated, OpenCL-C does not permit private pointers as kernel arguments and our intent was to keep maintain that restriction in SPIR-V binaries used for OpenCL.

Well, it's not technically a private pointer kernel argument, this is just an implementation detail on the compiler side (LLVM), because there is no other way of passing a struct parameter in LLVM. It's of course still handled as a normal value on the host side. And I would argue that SPIR-V has the exact same limitation and handling (see below).

For structs, this has the added benefit that the clSetKernelArg() validation is trivial for implementations. clSetKernelArg needs to validate the size of the argument being set. Passing a struct by private pointer in SPIR-V, but by value in the host API, opens further complications that the compiler must be able to determine the size of the struct, which could be non-trivial.

The goal of SPIR-V is to directly (when possible) represent the source program. The example source program has a struct in the parameter list. I would therefore expect OpTypeStruct to show up in the parameter list as well.

I was actually under the impression that the compiler already knows the size of the struct (which is fairly trivial to do, in my opinion), but it's awkwardly missing from the SPIR metadata that gets computed and emitted. Though SPIRV-LLVM will add the MaxByteOffset decoration/attribute if the parameter has a dereferenceable(#bytes) attribute on the LLVM side.

The issue with directly using OpTypeStruct (as already stated in the original post) is that it's impossible to do, because it would no longer be a pointer. If it's not a pointer, then it can't use Op * AccessChain (GEP) any more (and other things) and would be restricted to using OpCompositeExtract/OpCompositeInsert (extractvalue/insertvalue), which would in turn lose you the ability to dynamically access anything in the struct, because OpCompositeExtract/OpCompositeInsert can only have constant literal indices, compared to Op * AccessChain which allows dynamic index values. Thus, the parameter must be a pointer somehow.

Also, I don't see how there would be any difference between OpTypeStruct and a pointer to OpTypeStruct for host-side validation. The underlying types/data are the same.

[dneto0]

@a2flo wrote:

In this case, yes, it does add the byval attribute. But unless it is actually specified somewhere that this attribute is required, I'd consider all attributes to be optional and this should work regardless.

Maybe it's a picky point, but the SPIR-V spec doesn't say anything about LLVM. What the SPIRV-LLVM software project does is its own business, doesn't require a spec, and could change. Khronos only specifies what's valid SPIR-V. That might have holes and imperfections, sure, but it's not about LLVM.

Let's keep discussion of LLVM-related issues on the SPIRV-LLVM bug, and keep this bug for SPIR-V spec issues.

Also later:

I was actually under the impression that the compiler already knows the size of the struct (which is fairly trivial to do, in my opinion), but it's awkwardly missing from the SPIR metadata that gets computed and emitted.

Actually, that's not so straightforward. The size of a struct depends on its member types, and some basic types might differ in size between the host and device. That's why OpenCL 2.0 6.9 paragraph k bans kernel arguments of those troublesome types (and structs containing those troublesome types). And that doesn't even address alignment differences possible between host and device.

Lastly, I don't see why option 2 from your original post is "impossible". If you want to pass a struct value in SPIR-V, then just pass a struct value. (Ignoring whether OpenCL implementations should accept such a thing.) I don't see a requirement to get a pointer to the struct base. A struct is a value, not storage. ?
If the source language has a requirement that a function parameter also serves as a variable (storage) with that value, then you're really talking about having an OpFunctionParameter which is the value and then an OpVariable with Function storage class whose initializer is that OpFunctionParameter. And in that case it's squarely the compiler's responsibility to generate that SPIR-V code.

[raunintc]

Building on @dneto0 - we built SPIR-V to be able to be able to directly represent the source program. If the original source program passed a struct, then the SPIR-V should be passing a struct and the compiler.

Compilers are generally broken into 2 or 3 stages. Sizing a struct is the job of the first stage, the language front end, which may add padding and alignment adjustments to the final in-memory storage size of the struct. By the time the parsed language is sent down the compiler pipeline there's generally no difference between a user defined field and the padding/alignment data that may get appended. Sometime this matters to a program, sometimes not. Languages generally have a way to exert control over how a struct is laid out in memory for those cases where the programmer needs such control.

I'll push back a bit on the "impossible" claim as well. Yes, some technologies model byval struct passing as a pointer, others do not. SPIR-V will allow either method. SPIR-V also wants to represent the source program as close as possible, which leads me to the conclusion that OpTypeStruct should appear in the parameter list.

[sheredom]

This issue fell off the wagon slightly, but we (the SPIR-V group) believes the latest version of the OpenCL environment specification for SPIR-V (https://www.khronos.org/registry/OpenCL/specs/opencl-2.2-environment.pdf) fixes this issue.