llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: David Rivera (RiverDave) <details> <summary>Changes</summary> Follow up to the discussion/feedback received here: https://github.com/llvm/llvm-project/pull/196079#discussion_r3352210538 --- Full diff: https://github.com/llvm/llvm-project/pull/204663.diff 2 Files Affected: - (modified) clang/lib/CIR/CodeGen/CIRGenModule.cpp (+1-1) - (modified) clang/test/CIR/CodeGenCUDA/address-spaces.cu (+17-17) ``````````diff diff --git a/clang/lib/CIR/CodeGen/CIRGenModule.cpp b/clang/lib/CIR/CodeGen/CIRGenModule.cpp index b377f84e8d370..f240bdd146255 100644 --- a/clang/lib/CIR/CodeGen/CIRGenModule.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenModule.cpp @@ -1449,7 +1449,7 @@ void CIRGenModule::emitGlobalVarDefinition(const clang::VarDecl *vd, if (getLangOpts().CUDA && (isCUDASharedVar || isCUDAShadowVar || isCUDADeviceShadowVar)) { - init = cir::PoisonAttr::get(convertType(vd->getType())); + init = cir::UndefAttr::get(convertType(vd->getType())); } else if (vd->hasAttr<LoaderUninitializedAttr>()) { errorNYI(vd->getSourceRange(), "emitGlobalVarDefinition: loader uninitialized attribute"); diff --git a/clang/test/CIR/CodeGenCUDA/address-spaces.cu b/clang/test/CIR/CodeGenCUDA/address-spaces.cu index 2443384ef0649..9665ad330beae 100644 --- a/clang/test/CIR/CodeGenCUDA/address-spaces.cu +++ b/clang/test/CIR/CodeGenCUDA/address-spaces.cu @@ -52,8 +52,8 @@ // CIR-POST: cir.global external target_address_space(1) @i = #cir.int<0> // LLVM-DEVICE-DAG: @i = addrspace(1) {{.*}}global i32 0 // OGCG-DAG: @i = addrspace(1) externally_initialized global i32 0 -// CIR-HOST: cir.global {{.*}} @i = #cir.poison : {{.*}} {{{.*}}, cu.var_registration = #cir.cu.var_registration<i, Variable>} -// LLVM-HOST: @i = internal global i32 poison +// CIR-HOST: cir.global {{.*}} @i = #cir.undef : {{.*}} {{{.*}}, cu.var_registration = #cir.cu.var_registration<i, Variable>} +// LLVM-HOST: @i = internal global i32 undef // OGCG-HOST: @i = internal global i32 undef __device__ int i; @@ -62,24 +62,24 @@ __device__ int i; // CIR-DEVICE: cir.global constant external target_address_space(4) @j = #cir.int<0> : {{.*}} {{{.*}}, cu.externally_initialized = #cir.cu.externally_initialized, cu.var_registration = #cir.cu.var_registration<j, Variable, constant>} // LLVM-DEVICE-DAG: @j = addrspace(4) {{.*}}constant i32 0 // OGCG-DAG: @j = addrspace(4) externally_initialized constant i32 0 -// CIR-HOST: cir.global {{.*}} @j = #cir.poison : {{.*}} {{{.*}}, cu.var_registration = #cir.cu.var_registration<j, Variable, constant>} -// LLVM-HOST: @j = internal global i32 poison +// CIR-HOST: cir.global {{.*}} @j = #cir.undef : {{.*}} {{{.*}}, cu.var_registration = #cir.cu.var_registration<j, Variable, constant>} +// LLVM-HOST: @j = internal global i32 undef // OGCG-HOST: @j = internal global i32 undef __constant__ int j; -// CIR-PRE: cir.global external lang_address_space(offload_local) @k = #cir.poison -// CIR-POST: cir.global external target_address_space(3) @k = #cir.poison -// CIR-DEVICE: cir.global external target_address_space(3) @k = #cir.poison -// LLVM-DEVICE-DAG: @k = addrspace(3) global i32 {{undef|poison}} +// CIR-PRE: cir.global external lang_address_space(offload_local) @k = #cir.undef +// CIR-POST: cir.global external target_address_space(3) @k = #cir.undef +// CIR-DEVICE: cir.global external target_address_space(3) @k = #cir.undef +// LLVM-DEVICE-DAG: @k = addrspace(3) global i32 undef // OGCG-DAG: @k = addrspace(3) global i32 undef -// CIR-HOST: cir.global {{.*}} @k = #cir.poison -// LLVM-HOST: @k = internal global i32 poison +// CIR-HOST: cir.global {{.*}} @k = #cir.undef +// LLVM-HOST: @k = internal global i32 undef // OGCG-HOST: @k = internal global i32 undef __shared__ int k; -// CIR-PRE: cir.global external lang_address_space(offload_local) @b = #cir.poison : !cir.float -// CIR-POST: cir.global external target_address_space(3) @b = #cir.poison : !cir.float -// LLVM-DEVICE-DAG: @b = addrspace(3) global float {{undef|poison}} +// CIR-PRE: cir.global external lang_address_space(offload_local) @b = #cir.undef : !cir.float +// CIR-POST: cir.global external target_address_space(3) @b = #cir.undef : !cir.float +// LLVM-DEVICE-DAG: @b = addrspace(3) global float undef // OGCG-DAG: @b = addrspace(3) global float undef __shared__ float b; @@ -100,8 +100,8 @@ extern __constant__ int ext_constant_var; // External device variables with definitions should be internal on host // CIR-DEVICE: cir.global external target_address_space(1) @ext_device_var_def = #cir.int<1> // LLVM-DEVICE: @ext_device_var_def = addrspace(1) externally_initialized global i32 1, align 4 -// CIR-HOST: cir.global "private" internal {{.*}} @ext_device_var_def = #cir.poison : {{.*}} {{{.*}}, cu.var_registration = #cir.cu.var_registration<ext_device_var_def, Variable>} -// LLVM-HOST: @ext_device_var_def = internal global i32 poison, align 4 +// CIR-HOST: cir.global "private" internal {{.*}} @ext_device_var_def = #cir.undef : {{.*}} {{{.*}}, cu.var_registration = #cir.cu.var_registration<ext_device_var_def, Variable>} +// LLVM-HOST: @ext_device_var_def = internal global i32 undef, align 4 // OGCG-HOST: @ext_device_var_def = internal global i32 // OGCG-DEVICE: @ext_device_var_def = addrspace(1) externally_initialized global i32 1, align 4 extern __device__ int ext_device_var_def; @@ -110,8 +110,8 @@ __device__ int ext_device_var_def = 1; // CIR-DEVICE: cir.global constant external target_address_space(4) @ext_constant_var_def = #cir.int<2> // LLVM-DEVICE: @ext_constant_var_def = addrspace(4) externally_initialized constant i32 2, align 4 // OGCG-DEVICE: @ext_constant_var_def = addrspace(4) externally_initialized constant i32 2, align 4 -// CIR-HOST: cir.global "private" internal {{.*}} @ext_constant_var_def = #cir.poison : {{.*}} {{{.*}}, cu.var_registration = #cir.cu.var_registration<ext_constant_var_def, Variable, constant>} -// LLVM-HOST: @ext_constant_var_def = internal global i32 poison, align 4 +// CIR-HOST: cir.global "private" internal {{.*}} @ext_constant_var_def = #cir.undef : {{.*}} {{{.*}}, cu.var_registration = #cir.cu.var_registration<ext_constant_var_def, Variable, constant>} +// LLVM-HOST: @ext_constant_var_def = internal global i32 undef, align 4 // OGCG-HOST: @ext_constant_var_def = internal global i32 extern __constant__ int ext_constant_var_def; __constant__ int ext_constant_var_def = 2; `````````` </details> https://github.com/llvm/llvm-project/pull/204663 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
