vasilito/RedBear-OS
1
0
Fork 0
Code Issues Pull Requests Actions 12 Packages Projects Releases Wiki Activity
Files
fc98cf31d3dfb7db260fec07d4addce8383654cf
RedBear-OS/local/recipes/dev/libclc/source/clang/test/CodeGenCUDA/device-var-init.cu
T
vasilito cb424d7448 build: static patch-sanity linter (shift-left the malformed-patch class) 
verify-patch-sanity.py validates every active recipe .patch has internally-
consistent hunk line counts — catching the 'malformed patch at line N' failure
at commit/CI/preflight time instead of hours into a cook. This cycle hit that
class three times (qtwaylandscanner, sddm, xwayland), each only discovered when
cookbook tried to apply the patch.

Running it across the repo found 29 latent malformed patches (validated against
GNU patch: e.g. relibc/P3-sysv-ipc reproduces 'malformed patch at line 22').
They were harmless only because they sit in vendored recipes (baked, not re-
applied) — but would fail on any version-bump re-derivation. --fix recounts the
hunk headers (body untouched) and repaired all 29.

Wired into build-preflight.sh (Phase 1.0D) and redbear-ci.yml, with a unit test
(test-patch-sanity.sh). Skips archived/legacy trees and unvalidatable formats
(empty placeholders, bare-@@ git hunks).
2026-08-01 05:13:02 +03:00

308 lines
14 KiB
Plaintext
Raw Blame History

// REQUIRES: nvptx-registered-target
// REQUIRES: amdgpu-registered-target
// Make sure we don't allow dynamic initialization for device
// variables, but accept empty constructors allowed by CUDA.
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -fcuda-is-device -std=c++11 \
// RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefixes=DEVICE,NVPTX %s
// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -std=c++11 \
// RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefixes=HOST %s
// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -std=c++11 \
// RUN: -fno-threadsafe-statics -emit-llvm -o - %s | FileCheck -check-prefixes=DEVICE,AMDGCN %s
#ifdef __clang__
#include "Inputs/cuda.h"
#endif
// Use the types we share with Sema tests.
#include "Inputs/cuda-initializers.h"
__device__ int d_v;
// DEVICE: @d_v ={{.*}} addrspace(1) externally_initialized global i32 0,
// HOST: @d_v = internal global i32 undef,
__shared__ int s_v;
// DEVICE: @s_v ={{.*}} addrspace(3) global i32 undef,
// HOST: @s_v = internal global i32 undef,
__constant__ int c_v;
// DEVICE: addrspace(4) externally_initialized constant i32 0,
// HOST: @c_v = internal global i32 undef,
__device__ int d_v_i = 1;
// DEVICE: @d_v_i ={{.*}} addrspace(1) externally_initialized global i32 1,
// HOST: @d_v_i = internal global i32 undef,
// For `static` device variables, assume they won't be addressed from the host
// side.
static __device__ int d_s_v_i = 1;
// DEVICE: @_ZL7d_s_v_i = internal addrspace(1) global i32 1,
// Dummy function to keep static variables referenced.
__device__ int foo() {
return d_s_v_i;
}
// trivial constructor -- allowed
__device__ T d_t;
// DEVICE: @d_t ={{.*}} addrspace(1) externally_initialized global %struct.T zeroinitializer
// HOST: @d_t = internal global %struct.T undef,
__shared__ T s_t;
// DEVICE: @s_t ={{.*}} addrspace(3) global %struct.T undef,
// HOST: @s_t = internal global %struct.T undef,
__constant__ T c_t;
// DEVICE: @c_t ={{.*}} addrspace(4) externally_initialized constant %struct.T zeroinitializer,
// HOST: @c_t = internal global %struct.T undef,
__device__ T d_t_i = {2};
// DEVICE: @d_t_i ={{.*}} addrspace(1) externally_initialized global %struct.T { i32 2 },
// HOST: @d_t_i = internal global %struct.T undef,
__constant__ T c_t_i = {2};
// DEVICE: @c_t_i ={{.*}} addrspace(4) externally_initialized constant %struct.T { i32 2 },
// HOST: @c_t_i = internal global %struct.T undef,
// empty constructor
__device__ EC d_ec;
// DEVICE: @d_ec ={{.*}} addrspace(1) externally_initialized global %struct.EC zeroinitializer,
// HOST: @d_ec = internal global %struct.EC undef,
__shared__ EC s_ec;
// DEVICE: @s_ec ={{.*}} addrspace(3) global %struct.EC undef,
// HOST: @s_ec = internal global %struct.EC undef,
__constant__ EC c_ec;
// DEVICE: @c_ec ={{.*}} addrspace(4) externally_initialized constant %struct.EC zeroinitializer,
// HOST: @c_ec = internal global %struct.EC undef
// empty destructor
__device__ ED d_ed;
// DEVICE: @d_ed ={{.*}} addrspace(1) externally_initialized global %struct.ED zeroinitializer,
// HOST: @d_ed = internal global %struct.ED undef,
__shared__ ED s_ed;
// DEVICE: @s_ed ={{.*}} addrspace(3) global %struct.ED undef,
// HOST: @s_ed = internal global %struct.ED undef,
__constant__ ED c_ed;
// DEVICE: @c_ed ={{.*}} addrspace(4) externally_initialized constant %struct.ED zeroinitializer,
// HOST: @c_ed = internal global %struct.ED undef,
__device__ ECD d_ecd;
// DEVICE: @d_ecd ={{.*}} addrspace(1) externally_initialized global %struct.ECD zeroinitializer,
// HOST: @d_ecd = internal global %struct.ECD undef,
__shared__ ECD s_ecd;
// DEVICE: @s_ecd ={{.*}} addrspace(3) global %struct.ECD undef,
// HOST: @s_ecd = internal global %struct.ECD undef,
__constant__ ECD c_ecd;
// DEVICE: @c_ecd ={{.*}} addrspace(4) externally_initialized constant %struct.ECD zeroinitializer,
// HOST: @c_ecd = internal global %struct.ECD undef,
// empty templated constructor -- allowed with no arguments
__device__ ETC d_etc;
// DEVICE: @d_etc ={{.*}} addrspace(1) externally_initialized global %struct.ETC zeroinitializer,
// HOST: @d_etc = internal global %struct.ETC undef,
__shared__ ETC s_etc;
// DEVICE: @s_etc ={{.*}} addrspace(3) global %struct.ETC undef,
// HOST: @s_etc = internal global %struct.ETC undef,
__constant__ ETC c_etc;
// DEVICE: @c_etc ={{.*}} addrspace(4) externally_initialized constant %struct.ETC zeroinitializer,
// HOST: @c_etc = internal global %struct.ETC undef,
__device__ NCFS d_ncfs;
// DEVICE: @d_ncfs ={{.*}} addrspace(1) externally_initialized global %struct.NCFS { i32 3 }
// HOST: @d_ncfs = internal global %struct.NCFS undef,
__constant__ NCFS c_ncfs;
// DEVICE: @c_ncfs ={{.*}} addrspace(4) externally_initialized constant %struct.NCFS { i32 3 }
// HOST: @c_ncfs = internal global %struct.NCFS undef,
// Regular base class -- allowed
__device__ T_B_T d_t_b_t;
// DEVICE: @d_t_b_t ={{.*}} addrspace(1) externally_initialized global %struct.T_B_T zeroinitializer,
// HOST: @d_t_b_t = internal global %struct.T_B_T undef,
__shared__ T_B_T s_t_b_t;
// DEVICE: @s_t_b_t ={{.*}} addrspace(3) global %struct.T_B_T undef,
// HOST: @s_t_b_t = internal global %struct.T_B_T undef,
__constant__ T_B_T c_t_b_t;
// DEVICE: @c_t_b_t ={{.*}} addrspace(4) externally_initialized constant %struct.T_B_T zeroinitializer,
// HOST: @c_t_b_t = internal global %struct.T_B_T undef,
// Incapsulated object of allowed class -- allowed
__device__ T_F_T d_t_f_t;
// DEVICE: @d_t_f_t ={{.*}} addrspace(1) externally_initialized global %struct.T_F_T zeroinitializer,
// HOST: @d_t_f_t = internal global %struct.T_F_T undef,
__shared__ T_F_T s_t_f_t;
// DEVICE: @s_t_f_t ={{.*}} addrspace(3) global %struct.T_F_T undef,
// HOST: @s_t_f_t = internal global %struct.T_F_T undef,
__constant__ T_F_T c_t_f_t;
// DEVICE: @c_t_f_t ={{.*}} addrspace(4) externally_initialized constant %struct.T_F_T zeroinitializer,
// HOST: @c_t_f_t = internal global %struct.T_F_T undef,
// array of allowed objects -- allowed
__device__ T_FA_T d_t_fa_t;
// DEVICE: @d_t_fa_t ={{.*}} addrspace(1) externally_initialized global %struct.T_FA_T zeroinitializer,
// HOST: @d_t_fa_t = internal global %struct.T_FA_T undef,
__shared__ T_FA_T s_t_fa_t;
// DEVICE: @s_t_fa_t ={{.*}} addrspace(3) global %struct.T_FA_T undef,
// HOST: @s_t_fa_t = internal global %struct.T_FA_T undef,
__constant__ T_FA_T c_t_fa_t;
// DEVICE: @c_t_fa_t ={{.*}} addrspace(4) externally_initialized constant %struct.T_FA_T zeroinitializer,
// HOST: @c_t_fa_t = internal global %struct.T_FA_T undef,
// Calling empty base class initializer is OK
__device__ EC_I_EC d_ec_i_ec;
// DEVICE: @d_ec_i_ec ={{.*}} addrspace(1) externally_initialized global %struct.EC_I_EC zeroinitializer,
// HOST: @d_ec_i_ec = internal global %struct.EC_I_EC undef,
__shared__ EC_I_EC s_ec_i_ec;
// DEVICE: @s_ec_i_ec ={{.*}} addrspace(3) global %struct.EC_I_EC undef,
// HOST: @s_ec_i_ec = internal global %struct.EC_I_EC undef,
__constant__ EC_I_EC c_ec_i_ec;
// DEVICE: @c_ec_i_ec ={{.*}} addrspace(4) externally_initialized constant %struct.EC_I_EC zeroinitializer,
// HOST: @c_ec_i_ec = internal global %struct.EC_I_EC undef,
// DEVICE: @_ZZ2dfvE4s_ec = internal addrspace(3) global %struct.EC undef
// DEVICE: @_ZZ2dfvE5s_etc = internal addrspace(3) global %struct.ETC undef
// DEVICE: @_ZZ2dfvE11const_array = internal addrspace(4) constant [5 x i32] [i32 1, i32 2, i32 3, i32 4, i32 5]
// DEVICE: @_ZZ2dfvE9const_int = internal addrspace(4) constant i32 123
// We should not emit global initializers for device-side variables.
// DEVICE-NOT: @__cxx_global_var_init
// Make sure that initialization restrictions do not apply to local
// variables.
__device__ void df() {
// NVPTX: %[[ec:.*]] = alloca %struct.EC
// NVPTX: %[[ed:.*]] = alloca %struct.ED
// NVPTX: %[[ecd:.*]] = alloca %struct.ECD
// NVPTX: %[[etc:.*]] = alloca %struct.ETC
// NVPTX: %[[uc:.*]] = alloca %struct.UC
// NVPTX: %[[ud:.*]] = alloca %struct.UD
// NVPTX: %[[eci:.*]] = alloca %struct.ECI
// NVPTX: %[[nec:.*]] = alloca %struct.NEC
// NVPTX: %[[ned:.*]] = alloca %struct.NED
// NVPTX: %[[ncv:.*]] = alloca %struct.NCV
// NVPTX: %[[vd:.*]] = alloca %struct.VD
// NVPTX: %[[ncf:.*]] = alloca %struct.NCF
// NVPTX: %[[ncfs:.*]] = alloca %struct.NCFS
// NVPTX: %[[utc:.*]] = alloca %struct.UTC
// NVPTX: %[[netc:.*]] = alloca %struct.NETC
// NVPTX: %[[ec_i_ec:.*]] = alloca %struct.EC_I_EC
// NVPTX: %[[ec_i_ec1:.*]] = alloca %struct.EC_I_EC1
// NVPTX: %[[t_v_t:.*]] = alloca %struct.T_V_T
// NVPTX: %[[t_b_nec:.*]] = alloca %struct.T_B_NEC
// NVPTX: %[[t_f_nec:.*]] = alloca %struct.T_F_NEC
// NVPTX: %[[t_fa_nec:.*]] = alloca %struct.T_FA_NEC
// NVPTX: %[[t_b_ned:.*]] = alloca %struct.T_B_NED
// NVPTX: %[[t_f_ned:.*]] = alloca %struct.T_F_NED
// NVPTX: %[[t_fa_ned:.*]] = alloca %struct.T_FA_NED
// AMDGCN: %[[ec:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ec to ptr
// AMDGCN: %[[ed:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ed to ptr
// AMDGCN: %[[ecd:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ecd to ptr
// AMDGCN: %[[etc:.*]] ={{.*}} addrspacecast ptr addrspace(5) %etc to ptr
// AMDGCN: %[[uc:.*]] ={{.*}} addrspacecast ptr addrspace(5) %uc to ptr
// AMDGCN: %[[ud:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ud to ptr
// AMDGCN: %[[eci:.*]] ={{.*}} addrspacecast ptr addrspace(5) %eci to ptr
// AMDGCN: %[[nec:.*]] ={{.*}} addrspacecast ptr addrspace(5) %nec to ptr
// AMDGCN: %[[ned:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ned to ptr
// AMDGCN: %[[ncv:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ncv to ptr
// AMDGCN: %[[vd:.*]] ={{.*}} addrspacecast ptr addrspace(5) %vd to ptr
// AMDGCN: %[[ncf:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ncf to ptr
// AMDGCN: %[[ncfs:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ncfs to ptr
// AMDGCN: %[[utc:.*]] ={{.*}} addrspacecast ptr addrspace(5) %utc to ptr
// AMDGCN: %[[netc:.*]] ={{.*}} addrspacecast ptr addrspace(5) %netc to ptr
// AMDGCN: %[[ec_i_ec:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ec_i_ec to ptr
// AMDGCN: %[[ec_i_ec1:.*]] ={{.*}} addrspacecast ptr addrspace(5) %ec_i_ec1 to ptr
// AMDGCN: %[[t_v_t:.*]] ={{.*}} addrspacecast ptr addrspace(5) %t_v_t to ptr
// AMDGCN: %[[t_b_nec:.*]] ={{.*}} addrspacecast ptr addrspace(5) %t_b_nec to ptr
// AMDGCN: %[[t_f_nec:.*]] ={{.*}} addrspacecast ptr addrspace(5) %t_f_nec to ptr
// AMDGCN: %[[t_fa_nec:.*]] ={{.*}} addrspacecast ptr addrspace(5) %t_fa_nec to ptr
// AMDGCN: %[[t_b_ned:.*]] ={{.*}} addrspacecast ptr addrspace(5) %t_b_ned to ptr
// AMDGCN: %[[t_f_ned:.*]] ={{.*}} addrspacecast ptr addrspace(5) %t_f_ned to ptr
// AMDGCN: %[[t_fa_ned:.*]] ={{.*}} addrspacecast ptr addrspace(5) %t_fa_ned to ptr
T t;
// DEVICE-NOT: call
EC ec;
// DEVICE: call void @_ZN2ECC1Ev(ptr {{[^,]*}} %[[ec]])
ED ed;
// DEVICE-NOT: call
ECD ecd;
// DEVICE: call void @_ZN3ECDC1Ev(ptr {{[^,]*}} %[[ecd]])
ETC etc;
// DEVICE: call void @_ZN3ETCC1IJEEEDpT_(ptr {{[^,]*}} %[[etc]])
UC uc;
// undefined constructor -- not allowed
// DEVICE: call void @_ZN2UCC1Ev(ptr {{[^,]*}} %[[uc]])
UD ud;
// undefined destructor -- not allowed
// DEVICE-NOT: call
ECI eci;
// empty constructor w/ initializer list -- not allowed
// DEVICE: call void @_ZN3ECIC1Ev(ptr {{[^,]*}} %[[eci]])
NEC nec;
// non-empty constructor -- not allowed
// DEVICE: call void @_ZN3NECC1Ev(ptr {{[^,]*}} %[[nec]])
// non-empty destructor -- not allowed
NED ned;
// no-constructor, virtual method -- not allowed
// DEVICE: call void @_ZN3NCVC1Ev(ptr {{[^,]*}} %[[ncv]])
NCV ncv;
// DEVICE-NOT: call
VD vd;
// DEVICE: call void @_ZN2VDC1Ev(ptr {{[^,]*}} %[[vd]])
NCF ncf;
// DEVICE: call void @_ZN3NCFC1Ev(ptr {{[^,]*}} %[[ncf]])
NCFS ncfs;
// DEVICE: call void @_ZN4NCFSC1Ev(ptr {{[^,]*}} %[[ncfs]])
UTC utc;
// DEVICE: call void @_ZN3UTCC1IJEEEDpT_(ptr {{[^,]*}} %[[utc]])
NETC netc;
// DEVICE: call void @_ZN4NETCC1IJEEEDpT_(ptr {{[^,]*}} %[[netc]])
T_B_T t_b_t;
// DEVICE-NOT: call
T_F_T t_f_t;
// DEVICE-NOT: call
T_FA_T t_fa_t;
// DEVICE-NOT: call
EC_I_EC ec_i_ec;
// DEVICE: call void @_ZN7EC_I_ECC1Ev(ptr {{[^,]*}} %[[ec_i_ec]])
EC_I_EC1 ec_i_ec1;
// DEVICE: call void @_ZN8EC_I_EC1C1Ev(ptr {{[^,]*}} %[[ec_i_ec1]])
T_V_T t_v_t;
// DEVICE: call void @_ZN5T_V_TC1Ev(ptr {{[^,]*}} %[[t_v_t]])
T_B_NEC t_b_nec;
// DEVICE: call void @_ZN7T_B_NECC1Ev(ptr {{[^,]*}} %[[t_b_nec]])
T_F_NEC t_f_nec;
// DEVICE: call void @_ZN7T_F_NECC1Ev(ptr {{[^,]*}} %[[t_f_nec]])
T_FA_NEC t_fa_nec;
// DEVICE: call void @_ZN8T_FA_NECC1Ev(ptr {{[^,]*}} %[[t_fa_nec]])
T_B_NED t_b_ned;
// DEVICE-NOT: call
T_F_NED t_f_ned;
// DEVICE-NOT: call
T_FA_NED t_fa_ned;
// DEVICE-NOT: call
static __shared__ EC s_ec;
// DEVICE-NOT: call void @_ZN2ECC1Ev(ptr addrspacecast (ptr addrspace(3) @_ZZ2dfvE4s_ec to ptr))
static __shared__ ETC s_etc;
// DEVICE-NOT: call void @_ZN3ETCC1IJEEEDpT_(ptr addrspacecast (ptr addrspace(3) @_ZZ2dfvE5s_etc to ptr))
static const int const_array[] = {1, 2, 3, 4, 5};
static const int const_int = 123;
// anchor point separating constructors and destructors
df(); // DEVICE: call void @_Z2dfv()
// Verify that we only call non-empty destructors
// DEVICE-NEXT: call void @_ZN8T_FA_NEDD1Ev(ptr {{[^,]*}} %[[t_fa_ned]])
// DEVICE-NEXT: call void @_ZN7T_F_NEDD1Ev(ptr {{[^,]*}} %[[t_f_ned]])
// DEVICE-NEXT: call void @_ZN7T_B_NEDD1Ev(ptr {{[^,]*}} %[[t_b_ned]])
// DEVICE-NEXT: call void @_ZN2VDD1Ev(ptr {{[^,]*}} %[[vd]])
// DEVICE-NEXT: call void @_ZN3NEDD1Ev(ptr {{[^,]*}} %[[ned]])
// DEVICE-NEXT: call void @_ZN2UDD1Ev(ptr {{[^,]*}} %[[ud]])
// DEVICE-NEXT: call void @_ZN3ECDD1Ev(ptr {{[^,]*}} %[[ecd]])
// DEVICE-NEXT: call void @_ZN2EDD1Ev(ptr {{[^,]*}} %[[ed]])
// DEVICE-NEXT: ret void
}
// We should not emit global init function.
// DEVICE-NOT: @_GLOBAL__sub_I
Reference in New Issue View Git Blame Copy Permalink