Indexing an array member of a struct produces a GEP with a non-pointer result

[Zaneham]

Indexing an array member of a struct produces a GEP whose result type is the
element type rather than a pointer to it, and whose base is undef. The
backend catches it; the frontend does not.

struct S { int b; int a[4]; };
__global__ void k(int *out)
{
    struct S s;
    s.a[0] = 7;
    out[0] = s.a[0];
}

$ kath --rv-elf gep.cu -o gep.elf
rv_isel: GEP result must be a pointer

--ir shows why:

%1 = gep i32, undef, 0
%3 = gep ptr<global, i32>, %0, 0
%4 = gep i32, undef, 0

The middle one is a normal out[0] and looks right. The other two are the
s.a[0] accesses: result type i32 instead of ptr<i32>, base undef
instead of the alloca.

Scalar members are fine, including file-scope structs reached through a
pointer, so it is specific to array-typed members. Nothing backend-specific
either, the bad GEP is in the IR before any backend runs.

Three things it is not, having spent a while on each. Not memcpy, not
struct-by-value, and not the array decay in lower_expr's AST_MEMBER case.
collect_struct resolves a field's type without looking at the extent, so
int a[4] registers as a bare int, and bir_lower.c:3474 is the call site
that handles file-scope structs. But adding the extent there alone changes the
output by nothing at all, so something upstream of it is wrong too.

Turned up while wiring the soft-float runtime for Tensix. sfp_unpacked_t has
a _pad[] member and zeroing it walks straight into this.