1// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck -check-prefix=FAKE %s
2// RUN: %clang_cc1 %s -triple amdgcn-amd-amdhsa -emit-llvm -o - | FileCheck -check-prefix=AMDGCN %s
3
4typedef struct {
5    int i;
6    float f; // At non-zero offset.
7} ArrayStruct;
8
9__constant ArrayStruct constant_array_struct = { 0, 0.0f };
10
11typedef struct {
12    __constant float* constant_float_ptr;
13} ConstantArrayPointerStruct;
14
15// FAKE: %struct.ConstantArrayPointerStruct = type { float addrspace(2)* }
16// FAKE: addrspace(2) constant %struct.ConstantArrayPointerStruct { float addrspace(2)* bitcast (i8 addrspace(2)* getelementptr (i8, i8 addrspace(2)* bitcast (%struct.ArrayStruct addrspace(2)* @constant_array_struct to i8 addrspace(2)*), i64 4) to float addrspace(2)*) }
17// AMDGCN: %struct.ConstantArrayPointerStruct = type { float addrspace(4)* }
18// AMDGCN: addrspace(4) constant %struct.ConstantArrayPointerStruct { float addrspace(4)* bitcast (i8 addrspace(4)* getelementptr (i8, i8 addrspace(4)* bitcast (%struct.ArrayStruct addrspace(4)* @constant_array_struct to i8 addrspace(4)*), i64 4) to float addrspace(4)*) }
19// Bug  18567
20__constant ConstantArrayPointerStruct constant_array_pointer_struct = {
21    &constant_array_struct.f
22};
23
24__kernel void initializer_cast_is_valid_crash()
25{
26  unsigned char v512[64] = {
27      0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,
28      0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,
29      0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,
30      0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x00,0x02,0x00
31  };
32
33}
34