handling of kernel struct parameters

[a2flo]

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.

(will be cross-posting to https://github.com/KhronosGroup/SPIRV-Headers/issues since I think this is a spec bug that at the very least requires some explicit text that mentions the correct behavior)