Skip to content

Commit 6a9ad5f

Browse files
darkbuckdarkbuck
committed
[cuda][hip] Add CUDA builtin surface/texture reference support.
Summary: - Even though the bindless surface/texture interfaces are promoted, there are still code using surface/texture references. For example, [PR#26400](https://bugs.llvm.org/show_bug.cgi?id=26400) reports the compilation issue for code using `tex2D` with texture references. For better compatibility, this patch proposes the support of surface/texture references. - Due to the absent documentation and magic headers, it's believed that `nvcc` does use builtins for texture support. From the limited NVVM documentation[^nvvm] and NVPTX backend texture/surface related tests[^test], it's believed that surface/texture references are supported by replacing their reference types, which are annotated with `device_builtin_surface_type`/`device_builtin_texture_type`, with the corresponding handle-like object types, `cudaSurfaceObject_t` or `cudaTextureObject_t`, in the device-side compilation. On the host side, that global handle variables are registered and will be established and updated later when corresponding binding/unbinding APIs are called[^bind]. Surface/texture references are most like device global variables but represented in different types on the host and device sides. - In this patch, the following changes are proposed to support that behavior: + Refine `device_builtin_surface_type` and `device_builtin_texture_type` attributes to be applied on `Type` decl only to check whether a variable is of the surface/texture reference type. + Add hooks in code generation to replace that reference types with the correponding object types as well as all accesses to them. In particular, `nvvm.texsurf.handle.internal` should be used to load object handles from global reference variables[^texsurf] as well as metadata annotations. + Generate host-side registration with proper template argument parsing. --- [^nvvm]: https://docs.nvidia.com/cuda/pdf/NVVM_IR_Specification.pdf [^test]: https://raw.githubusercontent.com/llvm/llvm-project/master/llvm/test/CodeGen/NVPTX/tex-read-cuda.ll [^bind]: See section 3.2.11.1.2 ``Texture reference API` in [CUDA C Programming Guide](https://docs.nvidia.com/cuda/pdf/CUDA_C_Programming_Guide.pdf). [^texsurf]: According to NVVM IR, `nvvm.texsurf.handle` should be used. But, the current backend doesn't have that supported. We may revise that later. Reviewers: tra, rjmccall, yaxunl, a.sidorin Subscribers: cfe-commits Tags: #clang Differential Revision: https://reviews.llvm.org/D76365
1 parent bd12ecb commit 6a9ad5f

22 files changed

+703
-59
lines changed

‎clang/include/clang/AST/Type.h

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2111,6 +2111,11 @@ class alignas(8) Type : public ExtQualsTypeCommonBase {
21112111
/// than implicitly __strong.
21122112
bool isObjCARCImplicitlyUnretainedType() const;
21132113

2114+
/// Check if the type is the CUDA device builtin surface type.
2115+
bool isCUDADeviceBuiltinSurfaceType() const;
2116+
/// Check if the type is the CUDA device builtin texture type.
2117+
bool isCUDADeviceBuiltinTextureType() const;
2118+
21142119
/// Return the implicit lifetime for this type, which must not be dependent.
21152120
Qualifiers::ObjCLifetime getObjCARCImplicitLifetime() const;
21162121

‎clang/include/clang/Basic/Attr.td

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1064,16 +1064,20 @@ def CUDADeviceBuiltin : IgnoredAttr {
10641064
let LangOpts = [CUDA];
10651065
}
10661066

1067-
def CUDADeviceBuiltinSurfaceType : IgnoredAttr {
1067+
def CUDADeviceBuiltinSurfaceType : InheritableAttr {
10681068
let Spellings = [GNU<"device_builtin_surface_type">,
10691069
Declspec<"__device_builtin_surface_type__">];
10701070
let LangOpts = [CUDA];
1071+
let Subjects = SubjectList<[CXXRecord]>;
1072+
let Documentation = [CUDADeviceBuiltinSurfaceTypeDocs];
10711073
}
10721074

1073-
def CUDADeviceBuiltinTextureType : IgnoredAttr {
1075+
def CUDADeviceBuiltinTextureType : InheritableAttr {
10741076
let Spellings = [GNU<"device_builtin_texture_type">,
10751077
Declspec<"__device_builtin_texture_type__">];
10761078
let LangOpts = [CUDA];
1079+
let Subjects = SubjectList<[CXXRecord]>;
1080+
let Documentation = [CUDADeviceBuiltinTextureTypeDocs];
10771081
}
10781082

10791083
def CUDAGlobal : InheritableAttr {

‎clang/include/clang/Basic/AttrDocs.td

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4624,6 +4624,28 @@ the initializer on host side.
46244624
}];
46254625
}
46264626

4627+
def CUDADeviceBuiltinSurfaceTypeDocs : Documentation {
4628+
let Category = DocCatType;
4629+
let Content = [{
4630+
The ``device_builtin_surface_type`` attribute can be applied to a class
4631+
template when declaring the surface reference. A surface reference variable
4632+
could be accessed on the host side and, on the device side, might be translated
4633+
into an internal surface object, which is established through surface bind and
4634+
unbind runtime APIs.
4635+
}];
4636+
}
4637+
4638+
def CUDADeviceBuiltinTextureTypeDocs : Documentation {
4639+
let Category = DocCatType;
4640+
let Content = [{
4641+
The ``device_builtin_texture_type`` attribute can be applied to a class
4642+
template when declaring the texture reference. A texture reference variable
4643+
could be accessed on the host side and, on the device side, might be translated
4644+
into an internal texture object, which is established through texture bind and
4645+
unbind runtime APIs.
4646+
}];
4647+
}
4648+
46274649
def LifetimeOwnerDocs : Documentation {
46284650
let Category = DocCatDecl;
46294651
let Content = [{

‎clang/include/clang/Basic/DiagnosticSemaKinds.td

Lines changed: 23 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -7967,6 +7967,29 @@ def err_cuda_ovl_target : Error<
79677967
def note_cuda_ovl_candidate_target_mismatch : Note<
79687968
"candidate template ignored: target attributes do not match">;
79697969

7970+
def err_cuda_device_builtin_surftex_cls_template : Error<
7971+
"illegal device builtin %select{surface|texture}0 reference "
7972+
"class template %1 declared here">;
7973+
def note_cuda_device_builtin_surftex_cls_should_have_n_args : Note<
7974+
"%0 needs to have exactly %1 template parameters">;
7975+
def note_cuda_device_builtin_surftex_cls_should_have_match_arg : Note<
7976+
"the %select{1st|2nd|3rd}1 template parameter of %0 needs to be "
7977+
"%select{a type|an integer or enum value}2">;
7978+
7979+
def err_cuda_device_builtin_surftex_ref_decl : Error<
7980+
"illegal device builtin %select{surface|texture}0 reference "
7981+
"type %1 declared here">;
7982+
def note_cuda_device_builtin_surftex_should_be_template_class : Note<
7983+
"%0 needs to be instantiated from a class template with proper "
7984+
"template arguments">;
7985+
def note_cuda_device_builtin_surftex_should_have_n_args : Note<
7986+
"%0 needs to be instantiated from a class template with exactly "
7987+
"%1 template arguments">;
7988+
def note_cuda_device_builtin_surftex_should_have_match_arg : Note<
7989+
"%0 needs to be instantiated from a class template with the "
7990+
"%select{1st|2nd|3rd}1 template argument as "
7991+
"%select{a type|an integral value}2">;
7992+
79707993
def warn_non_pod_vararg_with_format_string : Warning<
79717994
"cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic "
79727995
"%select{function|block|method|constructor}2; expected type from format "

‎clang/lib/AST/Type.cpp

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4084,6 +4084,20 @@ bool Type::isCARCBridgableType() const {
40844084
return Pointee->isVoidType() || Pointee->isRecordType();
40854085
}
40864086

4087+
/// Check if the specified type is the CUDA device builtin surface type.
4088+
bool Type::isCUDADeviceBuiltinSurfaceType() const {
4089+
if (const auto *RT = getAs<RecordType>())
4090+
return RT->getDecl()->hasAttr<CUDADeviceBuiltinSurfaceTypeAttr>();
4091+
return false;
4092+
}
4093+
4094+
/// Check if the specified type is the CUDA device builtin texture type.
4095+
bool Type::isCUDADeviceBuiltinTextureType() const {
4096+
if (const auto *RT = getAs<RecordType>())
4097+
return RT->getDecl()->hasAttr<CUDADeviceBuiltinTextureTypeAttr>();
4098+
return false;
4099+
}
4100+
40874101
bool Type::hasSizedVLAType() const {
40884102
if (!isVariablyModifiedType()) return false;
40894103

‎clang/lib/CodeGen/CGCUDANV.cpp

Lines changed: 66 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -50,7 +50,7 @@ class CGNVCUDARuntime : public CGCUDARuntime {
5050
struct VarInfo {
5151
llvm::GlobalVariable *Var;
5252
const VarDecl *D;
53-
unsigned Flag;
53+
DeviceVarFlags Flags;
5454
};
5555
llvm::SmallVector<VarInfo, 16> DeviceVars;
5656
/// Keeps track of variable containing handle of GPU binary. Populated by
@@ -124,8 +124,25 @@ class CGNVCUDARuntime : public CGCUDARuntime {
124124

125125
void emitDeviceStub(CodeGenFunction &CGF, FunctionArgList &Args) override;
126126
void registerDeviceVar(const VarDecl *VD, llvm::GlobalVariable &Var,
127-
unsigned Flags) override {
128-
DeviceVars.push_back({&Var, VD, Flags});
127+
bool Extern, bool Constant) override {
128+
DeviceVars.push_back({&Var,
129+
VD,
130+
{DeviceVarFlags::Variable, Extern, Constant,
131+
/*Normalized*/ false, /*Type*/ 0}});
132+
}
133+
void registerDeviceSurf(const VarDecl *VD, llvm::GlobalVariable &Var,
134+
bool Extern, int Type) override {
135+
DeviceVars.push_back({&Var,
136+
VD,
137+
{DeviceVarFlags::Surface, Extern, /*Constant*/ false,
138+
/*Normalized*/ false, Type}});
139+
}
140+
void registerDeviceTex(const VarDecl *VD, llvm::GlobalVariable &Var,
141+
bool Extern, int Type, bool Normalized) override {
142+
DeviceVars.push_back({&Var,
143+
VD,
144+
{DeviceVarFlags::Texture, Extern, /*Constant*/ false,
145+
Normalized, Type}});
129146
}
130147

131148
/// Creates module constructor function
@@ -431,22 +448,55 @@ llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
431448
llvm::FunctionCallee RegisterVar = CGM.CreateRuntimeFunction(
432449
llvm::FunctionType::get(IntTy, RegisterVarParams, false),
433450
addUnderscoredPrefixToName("RegisterVar"));
451+
// void __cudaRegisterSurface(void **, const struct surfaceReference *,
452+
// const void **, const char *, int, int);
453+
llvm::FunctionCallee RegisterSurf = CGM.CreateRuntimeFunction(
454+
llvm::FunctionType::get(
455+
VoidTy, {VoidPtrPtrTy, VoidPtrTy, CharPtrTy, CharPtrTy, IntTy, IntTy},
456+
false),
457+
addUnderscoredPrefixToName("RegisterSurface"));
458+
// void __cudaRegisterTexture(void **, const struct textureReference *,
459+
// const void **, const char *, int, int, int)
460+
llvm::FunctionCallee RegisterTex = CGM.CreateRuntimeFunction(
461+
llvm::FunctionType::get(
462+
VoidTy,
463+
{VoidPtrPtrTy, VoidPtrTy, CharPtrTy, CharPtrTy, IntTy, IntTy, IntTy},
464+
false),
465+
addUnderscoredPrefixToName("RegisterTexture"));
434466
for (auto &&Info : DeviceVars) {
435467
llvm::GlobalVariable *Var = Info.Var;
436-
unsigned Flags = Info.Flag;
437468
llvm::Constant *VarName = makeConstantString(getDeviceSideName(Info.D));
438-
uint64_t VarSize =
439-
CGM.getDataLayout().getTypeAllocSize(Var->getValueType());
440-
llvm::Value *Args[] = {
441-
&GpuBinaryHandlePtr,
442-
Builder.CreateBitCast(Var, VoidPtrTy),
443-
VarName,
444-
VarName,
445-
llvm::ConstantInt::get(IntTy, (Flags & ExternDeviceVar) ? 1 : 0),
446-
llvm::ConstantInt::get(IntTy, VarSize),
447-
llvm::ConstantInt::get(IntTy, (Flags & ConstantDeviceVar) ? 1 : 0),
448-
llvm::ConstantInt::get(IntTy, 0)};
449-
Builder.CreateCall(RegisterVar, Args);
469+
switch (Info.Flags.Kind) {
470+
case DeviceVarFlags::Variable: {
471+
uint64_t VarSize =
472+
CGM.getDataLayout().getTypeAllocSize(Var->getValueType());
473+
llvm::Value *Args[] = {&GpuBinaryHandlePtr,
474+
Builder.CreateBitCast(Var, VoidPtrTy),
475+
VarName,
476+
VarName,
477+
llvm::ConstantInt::get(IntTy, Info.Flags.Extern),
478+
llvm::ConstantInt::get(IntTy, VarSize),
479+
llvm::ConstantInt::get(IntTy, Info.Flags.Constant),
480+
llvm::ConstantInt::get(IntTy, 0)};
481+
Builder.CreateCall(RegisterVar, Args);
482+
break;
483+
}
484+
case DeviceVarFlags::Surface:
485+
Builder.CreateCall(
486+
RegisterSurf,
487+
{&GpuBinaryHandlePtr, Builder.CreateBitCast(Var, VoidPtrTy), VarName,
488+
VarName, llvm::ConstantInt::get(IntTy, Info.Flags.SurfTexType),
489+
llvm::ConstantInt::get(IntTy, Info.Flags.Extern)});
490+
break;
491+
case DeviceVarFlags::Texture:
492+
Builder.CreateCall(
493+
RegisterTex,
494+
{&GpuBinaryHandlePtr, Builder.CreateBitCast(Var, VoidPtrTy), VarName,
495+
VarName, llvm::ConstantInt::get(IntTy, Info.Flags.SurfTexType),
496+
llvm::ConstantInt::get(IntTy, Info.Flags.Normalized),
497+
llvm::ConstantInt::get(IntTy, Info.Flags.Extern)});
498+
break;
499+
}
450500
}
451501

452502
Builder.CreateRetVoid();

‎clang/lib/CodeGen/CGCUDARuntime.h

Lines changed: 16 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -42,9 +42,17 @@ class CGCUDARuntime {
4242

4343
public:
4444
// Global variable properties that must be passed to CUDA runtime.
45-
enum DeviceVarFlags {
46-
ExternDeviceVar = 0x01, // extern
47-
ConstantDeviceVar = 0x02, // __constant__
45+
struct DeviceVarFlags {
46+
enum DeviceVarKind : unsigned {
47+
Variable, // Variable
48+
Surface, // Builtin surface
49+
Texture, // Builtin texture
50+
};
51+
DeviceVarKind Kind : 2;
52+
unsigned Extern : 1;
53+
unsigned Constant : 1; // Constant variable.
54+
unsigned Normalized : 1; // Normalized texture.
55+
int SurfTexType; // Type of surface/texutre.
4856
};
4957

5058
CGCUDARuntime(CodeGenModule &CGM) : CGM(CGM) {}
@@ -57,7 +65,11 @@ class CGCUDARuntime {
5765
/// Emits a kernel launch stub.
5866
virtual void emitDeviceStub(CodeGenFunction &CGF, FunctionArgList &Args) = 0;
5967
virtual void registerDeviceVar(const VarDecl *VD, llvm::GlobalVariable &Var,
60-
unsigned Flags) = 0;
68+
bool Extern, bool Constant) = 0;
69+
virtual void registerDeviceSurf(const VarDecl *VD, llvm::GlobalVariable &Var,
70+
bool Extern, int Type) = 0;
71+
virtual void registerDeviceTex(const VarDecl *VD, llvm::GlobalVariable &Var,
72+
bool Extern, int Type, bool Normalized) = 0;
6173

6274
/// Constructs and returns a module initialization function or nullptr if it's
6375
/// not needed. Must be called after all kernels have been emitted.

‎clang/lib/CodeGen/CGExprAgg.cpp

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -15,6 +15,7 @@
1515
#include "CodeGenFunction.h"
1616
#include "CodeGenModule.h"
1717
#include "ConstantEmitter.h"
18+
#include "TargetInfo.h"
1819
#include "clang/AST/ASTContext.h"
1920
#include "clang/AST/Attr.h"
2021
#include "clang/AST/DeclCXX.h"
@@ -1946,6 +1947,18 @@ void CodeGenFunction::EmitAggregateCopy(LValue Dest, LValue Src, QualType Ty,
19461947
}
19471948
}
19481949

1950+
if (getLangOpts().CUDAIsDevice) {
1951+
if (Ty->isCUDADeviceBuiltinSurfaceType()) {
1952+
if (getTargetHooks().emitCUDADeviceBuiltinSurfaceDeviceCopy(*this, Dest,
1953+
Src))
1954+
return;
1955+
} else if (Ty->isCUDADeviceBuiltinTextureType()) {
1956+
if (getTargetHooks().emitCUDADeviceBuiltinTextureDeviceCopy(*this, Dest,
1957+
Src))
1958+
return;
1959+
}
1960+
}
1961+
19491962
// Aggregate assignment turns into llvm.memcpy. This is almost valid per
19501963
// C99 6.5.16.1p3, which states "If the value being stored in an object is
19511964
// read from another object that overlaps in anyway the storage of the first

‎clang/lib/CodeGen/CodeGenModule.cpp

Lines changed: 57 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -713,6 +713,19 @@ llvm::MDNode *CodeGenModule::getTBAATypeInfo(QualType QTy) {
713713
TBAAAccessInfo CodeGenModule::getTBAAAccessInfo(QualType AccessType) {
714714
if (!TBAA)
715715
return TBAAAccessInfo();
716+
if (getLangOpts().CUDAIsDevice) {
717+
// As CUDA builtin surface/texture types are replaced, skip generating TBAA
718+
// access info.
719+
if (AccessType->isCUDADeviceBuiltinSurfaceType()) {
720+
if (getTargetCodeGenInfo().getCUDADeviceBuiltinSurfaceDeviceType() !=
721+
nullptr)
722+
return TBAAAccessInfo();
723+
} else if (AccessType->isCUDADeviceBuiltinTextureType()) {
724+
if (getTargetCodeGenInfo().getCUDADeviceBuiltinTextureDeviceType() !=
725+
nullptr)
726+
return TBAAAccessInfo();
727+
}
728+
}
716729
return TBAA->getAccessInfo(AccessType);
717730
}
718731

@@ -2507,7 +2520,9 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) {
25072520
!Global->hasAttr<CUDAGlobalAttr>() &&
25082521
!Global->hasAttr<CUDAConstantAttr>() &&
25092522
!Global->hasAttr<CUDASharedAttr>() &&
2510-
!(LangOpts.HIP && Global->hasAttr<HIPPinnedShadowAttr>()))
2523+
!(LangOpts.HIP && Global->hasAttr<HIPPinnedShadowAttr>()) &&
2524+
!Global->getType()->isCUDADeviceBuiltinSurfaceType() &&
2525+
!Global->getType()->isCUDADeviceBuiltinTextureType())
25112526
return;
25122527
} else {
25132528
// We need to emit host-side 'shadows' for all global
@@ -3960,12 +3975,16 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
39603975
!getLangOpts().CUDAIsDevice &&
39613976
(D->hasAttr<CUDAConstantAttr>() || D->hasAttr<CUDADeviceAttr>() ||
39623977
D->hasAttr<CUDASharedAttr>());
3978+
bool IsCUDADeviceShadowVar =
3979+
getLangOpts().CUDAIsDevice &&
3980+
(D->getType()->isCUDADeviceBuiltinSurfaceType() ||
3981+
D->getType()->isCUDADeviceBuiltinTextureType());
39633982
// HIP pinned shadow of initialized host-side global variables are also
39643983
// left undefined.
39653984
bool IsHIPPinnedShadowVar =
39663985
getLangOpts().CUDAIsDevice && D->hasAttr<HIPPinnedShadowAttr>();
3967-
if (getLangOpts().CUDA &&
3968-
(IsCUDASharedVar || IsCUDAShadowVar || IsHIPPinnedShadowVar))
3986+
if (getLangOpts().CUDA && (IsCUDASharedVar || IsCUDAShadowVar ||
3987+
IsCUDADeviceShadowVar || IsHIPPinnedShadowVar))
39693988
Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
39703989
else if (D->hasAttr<LoaderUninitializedAttr>())
39713990
Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
@@ -4076,25 +4095,48 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
40764095
if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
40774096
D->hasAttr<HIPPinnedShadowAttr>()) {
40784097
Linkage = llvm::GlobalValue::InternalLinkage;
4079-
4080-
// Shadow variables and their properties must be registered
4081-
// with CUDA runtime.
4082-
unsigned Flags = 0;
4083-
if (!D->hasDefinition())
4084-
Flags |= CGCUDARuntime::ExternDeviceVar;
4085-
if (D->hasAttr<CUDAConstantAttr>())
4086-
Flags |= CGCUDARuntime::ConstantDeviceVar;
4087-
// Extern global variables will be registered in the TU where they are
4088-
// defined.
4098+
// Shadow variables and their properties must be registered with CUDA
4099+
// runtime. Skip Extern global variables, which will be registered in
4100+
// the TU where they are defined.
40894101
if (!D->hasExternalStorage())
4090-
getCUDARuntime().registerDeviceVar(D, *GV, Flags);
4091-
} else if (D->hasAttr<CUDASharedAttr>())
4102+
getCUDARuntime().registerDeviceVar(D, *GV, !D->hasDefinition(),
4103+
D->hasAttr<CUDAConstantAttr>());
4104+
} else if (D->hasAttr<CUDASharedAttr>()) {
40924105
// __shared__ variables are odd. Shadows do get created, but
40934106
// they are not registered with the CUDA runtime, so they
40944107
// can't really be used to access their device-side
40954108
// counterparts. It's not clear yet whether it's nvcc's bug or
40964109
// a feature, but we've got to do the same for compatibility.
40974110
Linkage = llvm::GlobalValue::InternalLinkage;
4111+
} else if (D->getType()->isCUDADeviceBuiltinSurfaceType() ||
4112+
D->getType()->isCUDADeviceBuiltinTextureType()) {
4113+
// Builtin surfaces and textures and their template arguments are
4114+
// also registered with CUDA runtime.
4115+
Linkage = llvm::GlobalValue::InternalLinkage;
4116+
const ClassTemplateSpecializationDecl *TD =
4117+
cast<ClassTemplateSpecializationDecl>(
4118+
D->getType()->getAs<RecordType>()->getDecl());
4119+
const TemplateArgumentList &Args = TD->getTemplateInstantiationArgs();
4120+
if (TD->hasAttr<CUDADeviceBuiltinSurfaceTypeAttr>()) {
4121+
assert(Args.size() == 2 &&
4122+
"Unexpected number of template arguments of CUDA device "
4123+
"builtin surface type.");
4124+
auto SurfType = Args[1].getAsIntegral();
4125+
if (!D->hasExternalStorage())
4126+
getCUDARuntime().registerDeviceSurf(D, *GV, !D->hasDefinition(),
4127+
SurfType.getSExtValue());
4128+
} else {
4129+
assert(Args.size() == 3 &&
4130+
"Unexpected number of template arguments of CUDA device "
4131+
"builtin texture type.");
4132+
auto TexType = Args[1].getAsIntegral();
4133+
auto Normalized = Args[2].getAsIntegral();
4134+
if (!D->hasExternalStorage())
4135+
getCUDARuntime().registerDeviceTex(D, *GV, !D->hasDefinition(),
4136+
TexType.getSExtValue(),
4137+
Normalized.getZExtValue());
4138+
}
4139+
}
40984140
}
40994141
}
41004142

0 commit comments

Comments
 (0)