[cuda] Mark builtin texture/surface reference variable as 'externally_initialized'.

[hliao]

Summary

• They need to be preserved even if there's no reference within the device code as the host code may need to initialize them based on the application logic.

[yaxunl]

clang/lib/CodeGen/CodeGenModule.cpp

4441–4443

we need also do this with variables with HIPManagedAttr and add a test for that

[hliao]

clang/lib/CodeGen/CodeGenModule.cpp

4441–4443

those managed var seems already being marked 'externally_initalized' in lib/CodeGen/CGCUDANV.cpp. Do you have missing places to mark them?

[yaxunl]

LGTM on HIP side. I am not sure whether CUDA needs this.

clang/lib/CodeGen/CodeGenModule.cpp

4441–4443

sorry, I fogot. Then it should be OK. We don't have other places to handle managed variables.

[tra]

CUDA lowers texture/surface types into a special handle. I do not think externally_initialized matters for it.
AFAICT this change is a no-op for CUDA.

clang/lib/CodeGen/CodeGenModule.cpp

void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
   // CUDA B.2.2 "The __constant__ qualifier, optionally used together with
   // __device__, declares a variable that: [...]
   // Is accessible from all the threads within the grid and from the host
   // through the runtime library (cudaGetSymbolAddress() / cudaGetSymbolSize()
   // / cudaMemcpyToSymbol() / cudaMemcpyFromSymbol())."
   if (GV && LangOpts.CUDA) {
     if (LangOpts.CUDAIsDevice) {
       if (Linkage != llvm::GlobalValue::InternalLinkage &&
-          (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()))
+          (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
+           D->getType()->isCUDADeviceBuiltinSurfaceType() ||
+           D->getType()->isCUDADeviceBuiltinTextureType()))
         GV->setExternallyInitialized(true);
     } else {
       getCUDARuntime().internalizeDeviceSideVar(D, Linkage);
     }
     getCUDARuntime().handleVarRegistration(D, *GV);
   }
 
   GV->setInitializer(Init);

clang/test/CodeGenCUDA/surface.cu

 };
 
 // Partial specialization over `void`.
 template<int dim>
 struct __attribute__((device_builtin_surface_type)) surface<void, dim> : public surfaceReference {
 };
 
 // On the device side, surface references are represented as `i64` handles.
-// DEVICE: @surf ={{.*}} addrspace(1) global i64 undef, align 4
+// DEVICE: @surf ={{.*}} addrspace(1) externally_initialized global i64 undef, align 4
 // On the host side, they remain in the original type.
 // HOST: @surf = internal global %struct.surface
 // HOST: @0 = private unnamed_addr constant [5 x i8] c"surf\00"
 surface<void, 2> surf;
 
 __attribute__((device)) int suld_2d_zero(surface<void, 2>, int, int) asm("llvm.nvvm.suld.2d.i32.zero");
 
 // DEVICE-LABEL: i32 @_Z3fooii(i32 %x, i32 %y)

clang/test/CodeGenCUDA/texture.cu

enum ReadMode {
   NormalizedFloat = 1
 };
 
 template <typename T, int dim = 1, enum ReadMode mode = ElementType>
 struct __attribute__((device_builtin_texture_type)) texture : public textureReference {
 };
 
 // On the device side, texture references are represented as `i64` handles.
-// DEVICE: @tex ={{.*}} addrspace(1) global i64 undef, align 4
-// DEVICE: @norm ={{.*}} addrspace(1) global i64 undef, align 4
+// DEVICE: @tex ={{.*}} addrspace(1) externally_initialized global i64 undef, align 4
+// DEVICE: @norm ={{.*}} addrspace(1) externally_initialized global i64 undef, align 4
 // On the host side, they remain in the original type.
 // HOST: @tex = internal global %struct.texture
 // HOST: @norm = internal global %struct.texture
 // HOST: @0 = private unnamed_addr constant [4 x i8] c"tex\00"
 // HOST: @1 = private unnamed_addr constant [5 x i8] c"norm\00"
 texture<float, 2, ElementType> tex;
 texture<float, 2, NormalizedFloat> norm;