[llvm] fe8063e - Revert "[cuda][hip] Add CUDA builtin surface/texture reference support."
Artem Belevich via llvm-commits
llvm-commits at lists.llvm.org
Fri Mar 27 10:03:05 PDT 2020
Author: Artem Belevich
Date: 2020-03-27T10:01:38-07:00
New Revision: fe8063e1a0e983f1b4d38530f4fb157a26c0771c
URL: https://github.com/llvm/llvm-project/commit/fe8063e1a0e983f1b4d38530f4fb157a26c0771c
DIFF: https://github.com/llvm/llvm-project/commit/fe8063e1a0e983f1b4d38530f4fb157a26c0771c.diff
LOG: Revert "[cuda][hip] Add CUDA builtin surface/texture reference support."
This reverts commit 6a9ad5f3f4ac66f0cae592e911f4baeb6ee5eca6.
The patch breaks CUDA copmilation.
Differential Revision: https://reviews.llvm.org/D76365
Added:
Modified:
clang/include/clang/AST/Type.h
clang/include/clang/Basic/Attr.td
clang/include/clang/Basic/AttrDocs.td
clang/include/clang/Basic/DiagnosticSemaKinds.td
clang/lib/AST/Type.cpp
clang/lib/CodeGen/CGCUDANV.cpp
clang/lib/CodeGen/CGCUDARuntime.h
clang/lib/CodeGen/CGExprAgg.cpp
clang/lib/CodeGen/CodeGenModule.cpp
clang/lib/CodeGen/CodeGenTypes.cpp
clang/lib/CodeGen/TargetInfo.cpp
clang/lib/CodeGen/TargetInfo.h
clang/lib/Headers/__clang_cuda_runtime_wrapper.h
clang/lib/Sema/SemaDeclAttr.cpp
clang/lib/Sema/SemaDeclCXX.cpp
clang/test/Misc/pragma-attribute-supported-attributes-list.test
clang/test/SemaCUDA/attr-declspec.cu
clang/test/SemaCUDA/attributes-on-non-cuda.cu
clang/test/SemaCUDA/bad-attributes.cu
llvm/include/llvm/IR/Operator.h
Removed:
clang/test/CodeGenCUDA/surface.cu
clang/test/CodeGenCUDA/texture.cu
################################################################################
diff --git a/clang/include/clang/AST/Type.h b/clang/include/clang/AST/Type.h
index 6b46fc5ad312..3a2411b4ed29 100644
--- a/clang/include/clang/AST/Type.h
+++ b/clang/include/clang/AST/Type.h
@@ -2111,11 +2111,6 @@ class alignas(8) Type : public ExtQualsTypeCommonBase {
/// than implicitly __strong.
bool isObjCARCImplicitlyUnretainedType() const;
- /// Check if the type is the CUDA device builtin surface type.
- bool isCUDADeviceBuiltinSurfaceType() const;
- /// Check if the type is the CUDA device builtin texture type.
- bool isCUDADeviceBuiltinTextureType() const;
-
/// Return the implicit lifetime for this type, which must not be dependent.
Qualifiers::ObjCLifetime getObjCARCImplicitLifetime() const;
diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td
index 96bfdd313f47..5a90b2be2cbf 100644
--- a/clang/include/clang/Basic/Attr.td
+++ b/clang/include/clang/Basic/Attr.td
@@ -1064,20 +1064,16 @@ def CUDADeviceBuiltin : IgnoredAttr {
let LangOpts = [CUDA];
}
-def CUDADeviceBuiltinSurfaceType : InheritableAttr {
+def CUDADeviceBuiltinSurfaceType : IgnoredAttr {
let Spellings = [GNU<"device_builtin_surface_type">,
Declspec<"__device_builtin_surface_type__">];
let LangOpts = [CUDA];
- let Subjects = SubjectList<[CXXRecord]>;
- let Documentation = [CUDADeviceBuiltinSurfaceTypeDocs];
}
-def CUDADeviceBuiltinTextureType : InheritableAttr {
+def CUDADeviceBuiltinTextureType : IgnoredAttr {
let Spellings = [GNU<"device_builtin_texture_type">,
Declspec<"__device_builtin_texture_type__">];
let LangOpts = [CUDA];
- let Subjects = SubjectList<[CXXRecord]>;
- let Documentation = [CUDADeviceBuiltinTextureTypeDocs];
}
def CUDAGlobal : InheritableAttr {
diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td
index 2c89dc6f4952..a1cf25ed3929 100644
--- a/clang/include/clang/Basic/AttrDocs.td
+++ b/clang/include/clang/Basic/AttrDocs.td
@@ -4624,28 +4624,6 @@ the initializer on host side.
}];
}
-def CUDADeviceBuiltinSurfaceTypeDocs : Documentation {
- let Category = DocCatType;
- let Content = [{
-The ``device_builtin_surface_type`` attribute can be applied to a class
-template when declaring the surface reference. A surface reference variable
-could be accessed on the host side and, on the device side, might be translated
-into an internal surface object, which is established through surface bind and
-unbind runtime APIs.
- }];
-}
-
-def CUDADeviceBuiltinTextureTypeDocs : Documentation {
- let Category = DocCatType;
- let Content = [{
-The ``device_builtin_texture_type`` attribute can be applied to a class
-template when declaring the texture reference. A texture reference variable
-could be accessed on the host side and, on the device side, might be translated
-into an internal texture object, which is established through texture bind and
-unbind runtime APIs.
- }];
-}
-
def LifetimeOwnerDocs : Documentation {
let Category = DocCatDecl;
let Content = [{
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 044d35f19e20..8e26aa91a761 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -7967,29 +7967,6 @@ def err_cuda_ovl_target : Error<
def note_cuda_ovl_candidate_target_mismatch : Note<
"candidate template ignored: target attributes do not match">;
-def err_cuda_device_builtin_surftex_cls_template : Error<
- "illegal device builtin %select{surface|texture}0 reference "
- "class template %1 declared here">;
-def note_cuda_device_builtin_surftex_cls_should_have_n_args : Note<
- "%0 needs to have exactly %1 template parameters">;
-def note_cuda_device_builtin_surftex_cls_should_have_match_arg : Note<
- "the %select{1st|2nd|3rd}1 template parameter of %0 needs to be "
- "%select{a type|an integer or enum value}2">;
-
-def err_cuda_device_builtin_surftex_ref_decl : Error<
- "illegal device builtin %select{surface|texture}0 reference "
- "type %1 declared here">;
-def note_cuda_device_builtin_surftex_should_be_template_class : Note<
- "%0 needs to be instantiated from a class template with proper "
- "template arguments">;
-def note_cuda_device_builtin_surftex_should_have_n_args : Note<
- "%0 needs to be instantiated from a class template with exactly "
- "%1 template arguments">;
-def note_cuda_device_builtin_surftex_should_have_match_arg : Note<
- "%0 needs to be instantiated from a class template with the "
- "%select{1st|2nd|3rd}1 template argument as "
- "%select{a type|an integral value}2">;
-
def warn_non_pod_vararg_with_format_string : Warning<
"cannot pass %select{non-POD|non-trivial}0 object of type %1 to variadic "
"%select{function|block|method|constructor}2; expected type from format "
diff --git a/clang/lib/AST/Type.cpp b/clang/lib/AST/Type.cpp
index 9d4b77ec99bf..69c942e46f72 100644
--- a/clang/lib/AST/Type.cpp
+++ b/clang/lib/AST/Type.cpp
@@ -4084,20 +4084,6 @@ bool Type::isCARCBridgableType() const {
return Pointee->isVoidType() || Pointee->isRecordType();
}
-/// Check if the specified type is the CUDA device builtin surface type.
-bool Type::isCUDADeviceBuiltinSurfaceType() const {
- if (const auto *RT = getAs<RecordType>())
- return RT->getDecl()->hasAttr<CUDADeviceBuiltinSurfaceTypeAttr>();
- return false;
-}
-
-/// Check if the specified type is the CUDA device builtin texture type.
-bool Type::isCUDADeviceBuiltinTextureType() const {
- if (const auto *RT = getAs<RecordType>())
- return RT->getDecl()->hasAttr<CUDADeviceBuiltinTextureTypeAttr>();
- return false;
-}
-
bool Type::hasSizedVLAType() const {
if (!isVariablyModifiedType()) return false;
diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index ed02a7dc9173..5d8e545050d9 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -50,7 +50,7 @@ class CGNVCUDARuntime : public CGCUDARuntime {
struct VarInfo {
llvm::GlobalVariable *Var;
const VarDecl *D;
- DeviceVarFlags Flags;
+ unsigned Flag;
};
llvm::SmallVector<VarInfo, 16> DeviceVars;
/// Keeps track of variable containing handle of GPU binary. Populated by
@@ -124,25 +124,8 @@ class CGNVCUDARuntime : public CGCUDARuntime {
void emitDeviceStub(CodeGenFunction &CGF, FunctionArgList &Args) override;
void registerDeviceVar(const VarDecl *VD, llvm::GlobalVariable &Var,
- bool Extern, bool Constant) override {
- DeviceVars.push_back({&Var,
- VD,
- {DeviceVarFlags::Variable, Extern, Constant,
- /*Normalized*/ false, /*Type*/ 0}});
- }
- void registerDeviceSurf(const VarDecl *VD, llvm::GlobalVariable &Var,
- bool Extern, int Type) override {
- DeviceVars.push_back({&Var,
- VD,
- {DeviceVarFlags::Surface, Extern, /*Constant*/ false,
- /*Normalized*/ false, Type}});
- }
- void registerDeviceTex(const VarDecl *VD, llvm::GlobalVariable &Var,
- bool Extern, int Type, bool Normalized) override {
- DeviceVars.push_back({&Var,
- VD,
- {DeviceVarFlags::Texture, Extern, /*Constant*/ false,
- Normalized, Type}});
+ unsigned Flags) override {
+ DeviceVars.push_back({&Var, VD, Flags});
}
/// Creates module constructor function
@@ -448,55 +431,22 @@ llvm::Function *CGNVCUDARuntime::makeRegisterGlobalsFn() {
llvm::FunctionCallee RegisterVar = CGM.CreateRuntimeFunction(
llvm::FunctionType::get(IntTy, RegisterVarParams, false),
addUnderscoredPrefixToName("RegisterVar"));
- // void __cudaRegisterSurface(void **, const struct surfaceReference *,
- // const void **, const char *, int, int);
- llvm::FunctionCallee RegisterSurf = CGM.CreateRuntimeFunction(
- llvm::FunctionType::get(
- VoidTy, {VoidPtrPtrTy, VoidPtrTy, CharPtrTy, CharPtrTy, IntTy, IntTy},
- false),
- addUnderscoredPrefixToName("RegisterSurface"));
- // void __cudaRegisterTexture(void **, const struct textureReference *,
- // const void **, const char *, int, int, int)
- llvm::FunctionCallee RegisterTex = CGM.CreateRuntimeFunction(
- llvm::FunctionType::get(
- VoidTy,
- {VoidPtrPtrTy, VoidPtrTy, CharPtrTy, CharPtrTy, IntTy, IntTy, IntTy},
- false),
- addUnderscoredPrefixToName("RegisterTexture"));
for (auto &&Info : DeviceVars) {
llvm::GlobalVariable *Var = Info.Var;
+ unsigned Flags = Info.Flag;
llvm::Constant *VarName = makeConstantString(getDeviceSideName(Info.D));
- switch (Info.Flags.Kind) {
- case DeviceVarFlags::Variable: {
- uint64_t VarSize =
- CGM.getDataLayout().getTypeAllocSize(Var->getValueType());
- llvm::Value *Args[] = {&GpuBinaryHandlePtr,
- Builder.CreateBitCast(Var, VoidPtrTy),
- VarName,
- VarName,
- llvm::ConstantInt::get(IntTy, Info.Flags.Extern),
- llvm::ConstantInt::get(IntTy, VarSize),
- llvm::ConstantInt::get(IntTy, Info.Flags.Constant),
- llvm::ConstantInt::get(IntTy, 0)};
- Builder.CreateCall(RegisterVar, Args);
- break;
- }
- case DeviceVarFlags::Surface:
- Builder.CreateCall(
- RegisterSurf,
- {&GpuBinaryHandlePtr, Builder.CreateBitCast(Var, VoidPtrTy), VarName,
- VarName, llvm::ConstantInt::get(IntTy, Info.Flags.SurfTexType),
- llvm::ConstantInt::get(IntTy, Info.Flags.Extern)});
- break;
- case DeviceVarFlags::Texture:
- Builder.CreateCall(
- RegisterTex,
- {&GpuBinaryHandlePtr, Builder.CreateBitCast(Var, VoidPtrTy), VarName,
- VarName, llvm::ConstantInt::get(IntTy, Info.Flags.SurfTexType),
- llvm::ConstantInt::get(IntTy, Info.Flags.Normalized),
- llvm::ConstantInt::get(IntTy, Info.Flags.Extern)});
- break;
- }
+ uint64_t VarSize =
+ CGM.getDataLayout().getTypeAllocSize(Var->getValueType());
+ llvm::Value *Args[] = {
+ &GpuBinaryHandlePtr,
+ Builder.CreateBitCast(Var, VoidPtrTy),
+ VarName,
+ VarName,
+ llvm::ConstantInt::get(IntTy, (Flags & ExternDeviceVar) ? 1 : 0),
+ llvm::ConstantInt::get(IntTy, VarSize),
+ llvm::ConstantInt::get(IntTy, (Flags & ConstantDeviceVar) ? 1 : 0),
+ llvm::ConstantInt::get(IntTy, 0)};
+ Builder.CreateCall(RegisterVar, Args);
}
Builder.CreateRetVoid();
diff --git a/clang/lib/CodeGen/CGCUDARuntime.h b/clang/lib/CodeGen/CGCUDARuntime.h
index b26132420d65..330e950c98eb 100644
--- a/clang/lib/CodeGen/CGCUDARuntime.h
+++ b/clang/lib/CodeGen/CGCUDARuntime.h
@@ -42,17 +42,9 @@ class CGCUDARuntime {
public:
// Global variable properties that must be passed to CUDA runtime.
- struct DeviceVarFlags {
- enum DeviceVarKind : unsigned {
- Variable, // Variable
- Surface, // Builtin surface
- Texture, // Builtin texture
- };
- DeviceVarKind Kind : 2;
- unsigned Extern : 1;
- unsigned Constant : 1; // Constant variable.
- unsigned Normalized : 1; // Normalized texture.
- int SurfTexType; // Type of surface/texutre.
+ enum DeviceVarFlags {
+ ExternDeviceVar = 0x01, // extern
+ ConstantDeviceVar = 0x02, // __constant__
};
CGCUDARuntime(CodeGenModule &CGM) : CGM(CGM) {}
@@ -65,11 +57,7 @@ class CGCUDARuntime {
/// Emits a kernel launch stub.
virtual void emitDeviceStub(CodeGenFunction &CGF, FunctionArgList &Args) = 0;
virtual void registerDeviceVar(const VarDecl *VD, llvm::GlobalVariable &Var,
- bool Extern, bool Constant) = 0;
- virtual void registerDeviceSurf(const VarDecl *VD, llvm::GlobalVariable &Var,
- bool Extern, int Type) = 0;
- virtual void registerDeviceTex(const VarDecl *VD, llvm::GlobalVariable &Var,
- bool Extern, int Type, bool Normalized) = 0;
+ unsigned Flags) = 0;
/// Constructs and returns a module initialization function or nullptr if it's
/// not needed. Must be called after all kernels have been emitted.
diff --git a/clang/lib/CodeGen/CGExprAgg.cpp b/clang/lib/CodeGen/CGExprAgg.cpp
index fa2d228b7eeb..df576decd69d 100644
--- a/clang/lib/CodeGen/CGExprAgg.cpp
+++ b/clang/lib/CodeGen/CGExprAgg.cpp
@@ -15,7 +15,6 @@
#include "CodeGenFunction.h"
#include "CodeGenModule.h"
#include "ConstantEmitter.h"
-#include "TargetInfo.h"
#include "clang/AST/ASTContext.h"
#include "clang/AST/Attr.h"
#include "clang/AST/DeclCXX.h"
@@ -1947,18 +1946,6 @@ void CodeGenFunction::EmitAggregateCopy(LValue Dest, LValue Src, QualType Ty,
}
}
- if (getLangOpts().CUDAIsDevice) {
- if (Ty->isCUDADeviceBuiltinSurfaceType()) {
- if (getTargetHooks().emitCUDADeviceBuiltinSurfaceDeviceCopy(*this, Dest,
- Src))
- return;
- } else if (Ty->isCUDADeviceBuiltinTextureType()) {
- if (getTargetHooks().emitCUDADeviceBuiltinTextureDeviceCopy(*this, Dest,
- Src))
- return;
- }
- }
-
// Aggregate assignment turns into llvm.memcpy. This is almost valid per
// C99 6.5.16.1p3, which states "If the value being stored in an object is
// read from another object that overlaps in anyway the storage of the first
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp
index fbde1bf5cab7..b91c38ec9956 100644
--- a/clang/lib/CodeGen/CodeGenModule.cpp
+++ b/clang/lib/CodeGen/CodeGenModule.cpp
@@ -713,19 +713,6 @@ llvm::MDNode *CodeGenModule::getTBAATypeInfo(QualType QTy) {
TBAAAccessInfo CodeGenModule::getTBAAAccessInfo(QualType AccessType) {
if (!TBAA)
return TBAAAccessInfo();
- if (getLangOpts().CUDAIsDevice) {
- // As CUDA builtin surface/texture types are replaced, skip generating TBAA
- // access info.
- if (AccessType->isCUDADeviceBuiltinSurfaceType()) {
- if (getTargetCodeGenInfo().getCUDADeviceBuiltinSurfaceDeviceType() !=
- nullptr)
- return TBAAAccessInfo();
- } else if (AccessType->isCUDADeviceBuiltinTextureType()) {
- if (getTargetCodeGenInfo().getCUDADeviceBuiltinTextureDeviceType() !=
- nullptr)
- return TBAAAccessInfo();
- }
- }
return TBAA->getAccessInfo(AccessType);
}
@@ -2520,9 +2507,7 @@ void CodeGenModule::EmitGlobal(GlobalDecl GD) {
!Global->hasAttr<CUDAGlobalAttr>() &&
!Global->hasAttr<CUDAConstantAttr>() &&
!Global->hasAttr<CUDASharedAttr>() &&
- !(LangOpts.HIP && Global->hasAttr<HIPPinnedShadowAttr>()) &&
- !Global->getType()->isCUDADeviceBuiltinSurfaceType() &&
- !Global->getType()->isCUDADeviceBuiltinTextureType())
+ !(LangOpts.HIP && Global->hasAttr<HIPPinnedShadowAttr>()))
return;
} else {
// We need to emit host-side 'shadows' for all global
@@ -3922,16 +3907,12 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
!getLangOpts().CUDAIsDevice &&
(D->hasAttr<CUDAConstantAttr>() || D->hasAttr<CUDADeviceAttr>() ||
D->hasAttr<CUDASharedAttr>());
- bool IsCUDADeviceShadowVar =
- getLangOpts().CUDAIsDevice &&
- (D->getType()->isCUDADeviceBuiltinSurfaceType() ||
- D->getType()->isCUDADeviceBuiltinTextureType());
// HIP pinned shadow of initialized host-side global variables are also
// left undefined.
bool IsHIPPinnedShadowVar =
getLangOpts().CUDAIsDevice && D->hasAttr<HIPPinnedShadowAttr>();
- if (getLangOpts().CUDA && (IsCUDASharedVar || IsCUDAShadowVar ||
- IsCUDADeviceShadowVar || IsHIPPinnedShadowVar))
+ if (getLangOpts().CUDA &&
+ (IsCUDASharedVar || IsCUDAShadowVar || IsHIPPinnedShadowVar))
Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
else if (D->hasAttr<LoaderUninitializedAttr>())
Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
@@ -4042,48 +4023,25 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D,
if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
D->hasAttr<HIPPinnedShadowAttr>()) {
Linkage = llvm::GlobalValue::InternalLinkage;
- // Shadow variables and their properties must be registered with CUDA
- // runtime. Skip Extern global variables, which will be registered in
- // the TU where they are defined.
+
+ // Shadow variables and their properties must be registered
+ // with CUDA runtime.
+ unsigned Flags = 0;
+ if (!D->hasDefinition())
+ Flags |= CGCUDARuntime::ExternDeviceVar;
+ if (D->hasAttr<CUDAConstantAttr>())
+ Flags |= CGCUDARuntime::ConstantDeviceVar;
+ // Extern global variables will be registered in the TU where they are
+ // defined.
if (!D->hasExternalStorage())
- getCUDARuntime().registerDeviceVar(D, *GV, !D->hasDefinition(),
- D->hasAttr<CUDAConstantAttr>());
- } else if (D->hasAttr<CUDASharedAttr>()) {
+ getCUDARuntime().registerDeviceVar(D, *GV, Flags);
+ } else if (D->hasAttr<CUDASharedAttr>())
// __shared__ variables are odd. Shadows do get created, but
// they are not registered with the CUDA runtime, so they
// can't really be used to access their device-side
// counterparts. It's not clear yet whether it's nvcc's bug or
// a feature, but we've got to do the same for compatibility.
Linkage = llvm::GlobalValue::InternalLinkage;
- } else if (D->getType()->isCUDADeviceBuiltinSurfaceType() ||
- D->getType()->isCUDADeviceBuiltinTextureType()) {
- // Builtin surfaces and textures and their template arguments are
- // also registered with CUDA runtime.
- Linkage = llvm::GlobalValue::InternalLinkage;
- const ClassTemplateSpecializationDecl *TD =
- cast<ClassTemplateSpecializationDecl>(
- D->getType()->getAs<RecordType>()->getDecl());
- const TemplateArgumentList &Args = TD->getTemplateInstantiationArgs();
- if (TD->hasAttr<CUDADeviceBuiltinSurfaceTypeAttr>()) {
- assert(Args.size() == 2 &&
- "Unexpected number of template arguments of CUDA device "
- "builtin surface type.");
- auto SurfType = Args[1].getAsIntegral();
- if (!D->hasExternalStorage())
- getCUDARuntime().registerDeviceSurf(D, *GV, !D->hasDefinition(),
- SurfType.getSExtValue());
- } else {
- assert(Args.size() == 3 &&
- "Unexpected number of template arguments of CUDA device "
- "builtin texture type.");
- auto TexType = Args[1].getAsIntegral();
- auto Normalized = Args[2].getAsIntegral();
- if (!D->hasExternalStorage())
- getCUDARuntime().registerDeviceTex(D, *GV, !D->hasDefinition(),
- TexType.getSExtValue(),
- Normalized.getZExtValue());
- }
- }
}
}
diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp
index befd80de96f0..31eca16bbbe5 100644
--- a/clang/lib/CodeGen/CodeGenTypes.cpp
+++ b/clang/lib/CodeGen/CodeGenTypes.cpp
@@ -383,20 +383,6 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) {
const Type *Ty = T.getTypePtr();
- // For the device-side compilation, CUDA device builtin surface/texture types
- // may be represented in
diff erent types.
- if (Context.getLangOpts().CUDAIsDevice) {
- if (T->isCUDADeviceBuiltinSurfaceType()) {
- if (auto *Ty = CGM.getTargetCodeGenInfo()
- .getCUDADeviceBuiltinSurfaceDeviceType())
- return Ty;
- } else if (T->isCUDADeviceBuiltinTextureType()) {
- if (auto *Ty = CGM.getTargetCodeGenInfo()
- .getCUDADeviceBuiltinTextureDeviceType())
- return Ty;
- }
- }
-
// RecordTypes are cached and processed specially.
if (const RecordType *RT = dyn_cast<RecordType>(Ty))
return ConvertRecordDeclType(RT->getDecl());
diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp
index e64fe4f3943d..2b96cc40ae99 100644
--- a/clang/lib/CodeGen/TargetInfo.cpp
+++ b/clang/lib/CodeGen/TargetInfo.cpp
@@ -28,7 +28,6 @@
#include "llvm/ADT/Triple.h"
#include "llvm/ADT/Twine.h"
#include "llvm/IR/DataLayout.h"
-#include "llvm/IR/IntrinsicsNVPTX.h"
#include "llvm/IR/Type.h"
#include "llvm/Support/raw_ostream.h"
#include <algorithm> // std::sort
@@ -6415,14 +6414,9 @@ Address ARMABIInfo::EmitVAArg(CodeGenFunction &CGF, Address VAListAddr,
namespace {
-class NVPTXTargetCodeGenInfo;
-
class NVPTXABIInfo : public ABIInfo {
- NVPTXTargetCodeGenInfo &CGInfo;
-
public:
- NVPTXABIInfo(CodeGenTypes &CGT, NVPTXTargetCodeGenInfo &Info)
- : ABIInfo(CGT), CGInfo(Info) {}
+ NVPTXABIInfo(CodeGenTypes &CGT) : ABIInfo(CGT) {}
ABIArgInfo classifyReturnType(QualType RetTy) const;
ABIArgInfo classifyArgumentType(QualType Ty) const;
@@ -6435,61 +6429,16 @@ class NVPTXABIInfo : public ABIInfo {
class NVPTXTargetCodeGenInfo : public TargetCodeGenInfo {
public:
NVPTXTargetCodeGenInfo(CodeGenTypes &CGT)
- : TargetCodeGenInfo(new NVPTXABIInfo(CGT, *this)) {}
+ : TargetCodeGenInfo(new NVPTXABIInfo(CGT)) {}
void setTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
CodeGen::CodeGenModule &M) const override;
bool shouldEmitStaticExternCAliases() const override;
- llvm::Type *getCUDADeviceBuiltinSurfaceDeviceType() const override {
- // On the device side, surface reference is represented as an object handle
- // in 64-bit integer.
- return llvm::Type::getInt64Ty(getABIInfo().getVMContext());
- }
-
- llvm::Type *getCUDADeviceBuiltinTextureDeviceType() const override {
- // On the device side, texture reference is represented as an object handle
- // in 64-bit integer.
- return llvm::Type::getInt64Ty(getABIInfo().getVMContext());
- }
-
- bool emitCUDADeviceBuiltinSurfaceDeviceCopy(CodeGenFunction &CGF, LValue Dst,
- LValue Src) const override {
- emitBuiltinSurfTexDeviceCopy(CGF, Dst, Src);
- return true;
- }
-
- bool emitCUDADeviceBuiltinTextureDeviceCopy(CodeGenFunction &CGF, LValue Dst,
- LValue Src) const override {
- emitBuiltinSurfTexDeviceCopy(CGF, Dst, Src);
- return true;
- }
-
private:
- // Adds a NamedMDNode with GV, Name, and Operand as operands, and adds the
+ // Adds a NamedMDNode with F, Name, and Operand as operands, and adds the
// resulting MDNode to the nvvm.annotations MDNode.
- static void addNVVMMetadata(llvm::GlobalValue *GV, StringRef Name,
- int Operand);
-
- static void emitBuiltinSurfTexDeviceCopy(CodeGenFunction &CGF, LValue Dst,
- LValue Src) {
- llvm::Value *Handle = nullptr;
- llvm::Constant *C =
- llvm::dyn_cast<llvm::Constant>(Src.getAddress(CGF).getPointer());
- // Lookup `addrspacecast` through the constant pointer if any.
- if (auto *ASC = llvm::dyn_cast_or_null<llvm::AddrSpaceCastOperator>(C))
- C = llvm::cast<llvm::Constant>(ASC->getPointerOperand());
- if (auto *GV = llvm::dyn_cast_or_null<llvm::GlobalVariable>(C)) {
- // Load the handle from the specific global variable using
- // `nvvm.texsurf.handle.internal` intrinsic.
- Handle = CGF.EmitRuntimeCall(
- CGF.CGM.getIntrinsic(llvm::Intrinsic::nvvm_texsurf_handle_internal,
- {GV->getType()}),
- {GV}, "texsurf_handle");
- } else
- Handle = CGF.EmitLoadOfScalar(Src, SourceLocation());
- CGF.EmitStoreOfScalar(Handle, Dst);
- }
+ static void addNVVMMetadata(llvm::Function *F, StringRef Name, int Operand);
};
/// Checks if the type is unsupported directly by the current target.
@@ -6562,19 +6511,8 @@ ABIArgInfo NVPTXABIInfo::classifyArgumentType(QualType Ty) const {
Ty = EnumTy->getDecl()->getIntegerType();
// Return aggregates type as indirect by value
- if (isAggregateTypeForABI(Ty)) {
- // Under CUDA device compilation, tex/surf builtin types are replaced with
- // object types and passed directly.
- if (getContext().getLangOpts().CUDAIsDevice) {
- if (Ty->isCUDADeviceBuiltinSurfaceType())
- return ABIArgInfo::getDirect(
- CGInfo.getCUDADeviceBuiltinSurfaceDeviceType());
- if (Ty->isCUDADeviceBuiltinTextureType())
- return ABIArgInfo::getDirect(
- CGInfo.getCUDADeviceBuiltinTextureDeviceType());
- }
+ if (isAggregateTypeForABI(Ty))
return getNaturalAlignIndirect(Ty, /* byval */ true);
- }
return (Ty->isPromotableIntegerType() ? ABIArgInfo::getExtend(Ty)
: ABIArgInfo::getDirect());
@@ -6602,17 +6540,6 @@ void NVPTXTargetCodeGenInfo::setTargetAttributes(
const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const {
if (GV->isDeclaration())
return;
- const VarDecl *VD = dyn_cast_or_null<VarDecl>(D);
- if (VD) {
- if (M.getLangOpts().CUDA) {
- if (VD->getType()->isCUDADeviceBuiltinSurfaceType())
- addNVVMMetadata(GV, "surface", 1);
- else if (VD->getType()->isCUDADeviceBuiltinTextureType())
- addNVVMMetadata(GV, "texture", 1);
- return;
- }
- }
-
const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(D);
if (!FD) return;
@@ -6661,16 +6588,16 @@ void NVPTXTargetCodeGenInfo::setTargetAttributes(
}
}
-void NVPTXTargetCodeGenInfo::addNVVMMetadata(llvm::GlobalValue *GV,
- StringRef Name, int Operand) {
- llvm::Module *M = GV->getParent();
+void NVPTXTargetCodeGenInfo::addNVVMMetadata(llvm::Function *F, StringRef Name,
+ int Operand) {
+ llvm::Module *M = F->getParent();
llvm::LLVMContext &Ctx = M->getContext();
// Get "nvvm.annotations" metadata node
llvm::NamedMDNode *MD = M->getOrInsertNamedMetadata("nvvm.annotations");
llvm::Metadata *MDVals[] = {
- llvm::ConstantAsMetadata::get(GV), llvm::MDString::get(Ctx, Name),
+ llvm::ConstantAsMetadata::get(F), llvm::MDString::get(Ctx, Name),
llvm::ConstantAsMetadata::get(
llvm::ConstantInt::get(llvm::Type::getInt32Ty(Ctx), Operand))};
// Append metadata to nvvm.annotations
diff --git a/clang/lib/CodeGen/TargetInfo.h b/clang/lib/CodeGen/TargetInfo.h
index e7c842bae4a9..e1e90e73cb58 100644
--- a/clang/lib/CodeGen/TargetInfo.h
+++ b/clang/lib/CodeGen/TargetInfo.h
@@ -315,32 +315,6 @@ class TargetCodeGenInfo {
virtual bool shouldEmitStaticExternCAliases() const { return true; }
virtual void setCUDAKernelCallingConvention(const FunctionType *&FT) const {}
-
- /// Return the device-side type for the CUDA device builtin surface type.
- virtual llvm::Type *getCUDADeviceBuiltinSurfaceDeviceType() const {
- // By default, no change from the original one.
- return nullptr;
- }
- /// Return the device-side type for the CUDA device builtin texture type.
- virtual llvm::Type *getCUDADeviceBuiltinTextureDeviceType() const {
- // By default, no change from the original one.
- return nullptr;
- }
-
- /// Emit the device-side copy of the builtin surface type.
- virtual bool emitCUDADeviceBuiltinSurfaceDeviceCopy(CodeGenFunction &CGF,
- LValue Dst,
- LValue Src) const {
- // DO NOTHING by default.
- return false;
- }
- /// Emit the device-side copy of the builtin texture type.
- virtual bool emitCUDADeviceBuiltinTextureDeviceCopy(CodeGenFunction &CGF,
- LValue Dst,
- LValue Src) const {
- // DO NOTHING by default.
- return false;
- }
};
} // namespace CodeGen
diff --git a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h
index 349a4c759bca..e91de3c81dbd 100644
--- a/clang/lib/Headers/__clang_cuda_runtime_wrapper.h
+++ b/clang/lib/Headers/__clang_cuda_runtime_wrapper.h
@@ -83,15 +83,13 @@
#if CUDA_VERSION < 9000
#define __CUDABE__
#else
-#define __CUDACC__
#define __CUDA_LIBDEVICE__
#endif
// Disables definitions of device-side runtime support stubs in
// cuda_device_runtime_api.h
-#include "host_defines.h"
-#undef __CUDACC__
#include "driver_types.h"
#include "host_config.h"
+#include "host_defines.h"
// Temporarily replace "nv_weak" with weak, so __attribute__((nv_weak)) in
// cuda_device_runtime_api.h ends up being __attribute__((weak)) which is the
diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp
index 5aacb2fd64a2..061a7d0225ed 100644
--- a/clang/lib/Sema/SemaDeclAttr.cpp
+++ b/clang/lib/Sema/SemaDeclAttr.cpp
@@ -6934,16 +6934,6 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
handleSimpleAttributeWithExclusions<HIPPinnedShadowAttr, CUDADeviceAttr,
CUDAConstantAttr>(S, D, AL);
break;
- case ParsedAttr::AT_CUDADeviceBuiltinSurfaceType:
- handleSimpleAttributeWithExclusions<CUDADeviceBuiltinSurfaceTypeAttr,
- CUDADeviceBuiltinTextureTypeAttr>(S, D,
- AL);
- break;
- case ParsedAttr::AT_CUDADeviceBuiltinTextureType:
- handleSimpleAttributeWithExclusions<CUDADeviceBuiltinTextureTypeAttr,
- CUDADeviceBuiltinSurfaceTypeAttr>(S, D,
- AL);
- break;
case ParsedAttr::AT_GNUInline:
handleGNUInlineAttr(S, D, AL);
break;
diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp
index 3b121ae8df4d..b65dc5c6427b 100644
--- a/clang/lib/Sema/SemaDeclCXX.cpp
+++ b/clang/lib/Sema/SemaDeclCXX.cpp
@@ -5877,183 +5877,6 @@ static void checkForMultipleExportedDefaultConstructors(Sema &S,
}
}
-static void checkCUDADeviceBuiltinSurfaceClassTemplate(Sema &S,
- CXXRecordDecl *Class) {
- bool ErrorReported = false;
- auto reportIllegalClassTemplate = [&ErrorReported](Sema &S,
- CXXRecordDecl *RD) {
- if (ErrorReported)
- return;
- S.Diag(RD->getLocation(),
- diag::err_cuda_device_builtin_surftex_cls_template)
- << /*surface*/ 0 << RD;
- ErrorReported = true;
- };
-
- TemplateParameterList *Params =
- Class->getDescribedClassTemplate()->getTemplateParameters();
- unsigned N = Params->size();
-
- if (N != 2) {
- reportIllegalClassTemplate(S, Class);
- S.Diag(Class->getLocation(),
- diag::note_cuda_device_builtin_surftex_cls_should_have_n_args)
- << Class << 2;
- }
- if (N > 0 && !isa<TemplateTypeParmDecl>(Params->getParam(0))) {
- reportIllegalClassTemplate(S, Class);
- S.Diag(Class->getLocation(),
- diag::note_cuda_device_builtin_surftex_cls_should_have_match_arg)
- << Class << /*1st*/ 0 << /*type*/ 0;
- }
- if (N > 1) {
- auto *NTTP = dyn_cast<NonTypeTemplateParmDecl>(Params->getParam(1));
- if (!NTTP || !NTTP->getType()->isIntegralOrEnumerationType()) {
- reportIllegalClassTemplate(S, Class);
- S.Diag(Class->getLocation(),
- diag::note_cuda_device_builtin_surftex_cls_should_have_match_arg)
- << Class << /*2nd*/ 1 << /*integer*/ 1;
- }
- }
-}
-
-static void checkCUDADeviceBuiltinTextureClassTemplate(Sema &S,
- CXXRecordDecl *Class) {
- bool ErrorReported = false;
- auto reportIllegalClassTemplate = [&ErrorReported](Sema &S,
- CXXRecordDecl *RD) {
- if (ErrorReported)
- return;
- S.Diag(RD->getLocation(),
- diag::err_cuda_device_builtin_surftex_cls_template)
- << /*texture*/ 1 << RD;
- ErrorReported = true;
- };
-
- TemplateParameterList *Params =
- Class->getDescribedClassTemplate()->getTemplateParameters();
- unsigned N = Params->size();
-
- if (N != 3) {
- reportIllegalClassTemplate(S, Class);
- S.Diag(Class->getLocation(),
- diag::note_cuda_device_builtin_surftex_cls_should_have_n_args)
- << Class << 3;
- }
- if (N > 0 && !isa<TemplateTypeParmDecl>(Params->getParam(0))) {
- reportIllegalClassTemplate(S, Class);
- S.Diag(Class->getLocation(),
- diag::note_cuda_device_builtin_surftex_cls_should_have_match_arg)
- << Class << /*1st*/ 0 << /*type*/ 0;
- }
- if (N > 1) {
- auto *NTTP = dyn_cast<NonTypeTemplateParmDecl>(Params->getParam(1));
- if (!NTTP || !NTTP->getType()->isIntegralOrEnumerationType()) {
- reportIllegalClassTemplate(S, Class);
- S.Diag(Class->getLocation(),
- diag::note_cuda_device_builtin_surftex_cls_should_have_match_arg)
- << Class << /*2nd*/ 1 << /*integer*/ 1;
- }
- }
- if (N > 2) {
- auto *NTTP = dyn_cast<NonTypeTemplateParmDecl>(Params->getParam(2));
- if (!NTTP || !NTTP->getType()->isIntegralOrEnumerationType()) {
- reportIllegalClassTemplate(S, Class);
- S.Diag(Class->getLocation(),
- diag::note_cuda_device_builtin_surftex_cls_should_have_match_arg)
- << Class << /*3rd*/ 2 << /*integer*/ 1;
- }
- }
-}
-
-static void checkCUDADeviceBuiltinSurfaceType(Sema &S, CXXRecordDecl *Class) {
- bool ErrorReported = false;
- auto reportIllegalReferenceType = [&ErrorReported](Sema &S,
- CXXRecordDecl *RD) {
- if (ErrorReported)
- return;
- S.Diag(RD->getLocation(), diag::err_cuda_device_builtin_surftex_ref_decl)
- << /*surface*/ 0 << RD;
- ErrorReported = true;
- };
-
- const auto *TD = dyn_cast<ClassTemplateSpecializationDecl>(Class);
- if (!TD) {
- reportIllegalReferenceType(S, Class);
- S.Diag(Class->getLocation(),
- diag::note_cuda_device_builtin_surftex_should_be_template_class)
- << Class;
- return;
- }
- const auto &Args = TD->getTemplateInstantiationArgs();
- unsigned N = Args.size();
- if (N != 2) {
- reportIllegalReferenceType(S, Class);
- S.Diag(Class->getLocation(),
- diag::note_cuda_device_builtin_surftex_should_have_n_args)
- << Class << /*nargs*/ 2;
- }
- if (N > 0 && Args[0].getKind() != TemplateArgument::Type) {
- reportIllegalReferenceType(S, Class);
- S.Diag(Class->getLocation(),
- diag::note_cuda_device_builtin_surftex_should_have_match_arg)
- << Class << /*1st*/ 0 << /*type*/ 0;
- }
- if (N > 1 && Args[1].getKind() != TemplateArgument::Integral) {
- reportIllegalReferenceType(S, Class);
- S.Diag(Class->getLocation(),
- diag::note_cuda_device_builtin_surftex_should_have_match_arg)
- << Class << /*2nd*/ 1 << /*integral*/ 1;
- }
-}
-
-static void checkCUDADeviceBuiltinTextureType(Sema &S, CXXRecordDecl *Class) {
- bool ErrorReported = false;
- auto reportIllegalReferenceType = [&ErrorReported](Sema &S,
- CXXRecordDecl *RD) {
- if (ErrorReported)
- return;
- S.Diag(RD->getLocation(), diag::err_cuda_device_builtin_surftex_ref_decl)
- << /*texture*/ 1 << RD;
- ErrorReported = true;
- };
-
- const auto *TD = dyn_cast<ClassTemplateSpecializationDecl>(Class);
- if (!TD) {
- reportIllegalReferenceType(S, Class);
- S.Diag(Class->getLocation(),
- diag::note_cuda_device_builtin_surftex_should_be_template_class)
- << Class;
- return;
- }
- const auto &Args = TD->getTemplateInstantiationArgs();
- unsigned N = Args.size();
- if (N != 3) {
- reportIllegalReferenceType(S, Class);
- S.Diag(Class->getLocation(),
- diag::note_cuda_device_builtin_surftex_should_have_n_args)
- << Class << /*nargs*/ 3;
- }
- if (N > 0 && Args[0].getKind() != TemplateArgument::Type) {
- reportIllegalReferenceType(S, Class);
- S.Diag(Class->getLocation(),
- diag::note_cuda_device_builtin_surftex_should_have_match_arg)
- << Class << /*1st*/ 0 << /*type*/ 0;
- }
- if (N > 1 && Args[1].getKind() != TemplateArgument::Integral) {
- reportIllegalReferenceType(S, Class);
- S.Diag(Class->getLocation(),
- diag::note_cuda_device_builtin_surftex_should_have_match_arg)
- << Class << /*2nd*/ 1 << /*integral*/ 1;
- }
- if (N > 2 && Args[2].getKind() != TemplateArgument::Integral) {
- reportIllegalReferenceType(S, Class);
- S.Diag(Class->getLocation(),
- diag::note_cuda_device_builtin_surftex_should_have_match_arg)
- << Class << /*3rd*/ 2 << /*integral*/ 1;
- }
-}
-
void Sema::checkClassLevelCodeSegAttribute(CXXRecordDecl *Class) {
// Mark any compiler-generated routines with the implicit code_seg attribute.
for (auto *Method : Class->methods()) {
@@ -6834,20 +6657,6 @@ void Sema::CheckCompletedCXXClass(Scope *S, CXXRecordDecl *Record) {
// is especially required for cases like vtable assumption loads.
MarkVTableUsed(Record->getInnerLocStart(), Record);
}
-
- if (getLangOpts().CUDA) {
- if (Record->getDescribedClassTemplate()) {
- if (Record->hasAttr<CUDADeviceBuiltinSurfaceTypeAttr>())
- checkCUDADeviceBuiltinSurfaceClassTemplate(*this, Record);
- else if (Record->hasAttr<CUDADeviceBuiltinTextureTypeAttr>())
- checkCUDADeviceBuiltinTextureClassTemplate(*this, Record);
- } else {
- if (Record->hasAttr<CUDADeviceBuiltinSurfaceTypeAttr>())
- checkCUDADeviceBuiltinSurfaceType(*this, Record);
- else if (Record->hasAttr<CUDADeviceBuiltinTextureTypeAttr>())
- checkCUDADeviceBuiltinTextureType(*this, Record);
- }
- }
}
/// Look up the special member function that would be called by a special
diff --git a/clang/test/CodeGenCUDA/surface.cu b/clang/test/CodeGenCUDA/surface.cu
deleted file mode 100644
index c4c0a59f495b..000000000000
--- a/clang/test/CodeGenCUDA/surface.cu
+++ /dev/null
@@ -1,37 +0,0 @@
-// REQUIRES: x86-registered-target
-// REQUIRES: nvptx-registered-target
-
-// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -triple nvptx64-nvidia-cuda -emit-llvm -o - %s | FileCheck --check-prefix=DEVICE %s
-// RUN: echo "GPU binary would be here" > %t
-// RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu -target-sdk-version=8.0 -fcuda-include-gpubinary %t -emit-llvm -o - %s | FileCheck --check-prefix=HOST %s
-
-struct surfaceReference {
- int desc;
-};
-
-template <typename T, int type = 1>
-struct __attribute__((device_builtin_surface_type)) surface : public surfaceReference {
-};
-
-// On the device side, surface references are represented as `i64` handles.
-// DEVICE: @surf = addrspace(1) 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)
-// DEVICE: call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @surf)
-// DEVICE: call i32 @llvm.nvvm.suld.2d.i32.zero(i64 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})
-__attribute__((device)) int foo(int x, int y) {
- return suld_2d_zero(surf, x, y);
-}
-
-// HOST: define internal void @[[PREFIX:__cuda]]_register_globals
-// Texture references need registering with correct arguments.
-// HOST: call void @[[PREFIX]]RegisterSurface(i8** %0, i8*{{.*}}({{.*}}@surf{{.*}}), i8*{{.*}}({{.*}}@0{{.*}}), i8*{{.*}}({{.*}}@0{{.*}}), i32 2, i32 0)
-
-// They also need annotating in metadata.
-// DEVICE: !0 = !{i64 addrspace(1)* @surf, !"surface", i32 1}
diff --git a/clang/test/CodeGenCUDA/texture.cu b/clang/test/CodeGenCUDA/texture.cu
deleted file mode 100644
index 7838eeb110b0..000000000000
--- a/clang/test/CodeGenCUDA/texture.cu
+++ /dev/null
@@ -1,55 +0,0 @@
-// REQUIRES: x86-registered-target
-// REQUIRES: nvptx-registered-target
-
-// RUN: %clang_cc1 -std=c++11 -fcuda-is-device -triple nvptx64-nvidia-cuda -emit-llvm -o - %s | FileCheck --check-prefix=DEVICE %s
-// RUN: echo "GPU binary would be here" > %t
-// RUN: %clang_cc1 -std=c++11 -triple x86_64-unknown-linux-gnu -target-sdk-version=8.0 -fcuda-include-gpubinary %t -emit-llvm -o - %s | FileCheck --check-prefix=HOST %s
-
-struct textureReference {
- int desc;
-};
-
-enum ReadMode {
- ElementType = 0,
- 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
-// 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;
-
-struct v4f {
- float x, y, z, w;
-};
-
-__attribute__((device)) v4f tex2d_ld(texture<float, 2, ElementType>, float, float) asm("llvm.nvvm.tex.unified.2d.v4f32.f32");
-__attribute__((device)) v4f tex2d_ld(texture<float, 2, NormalizedFloat>, int, int) asm("llvm.nvvm.tex.unified.2d.v4f32.s32");
-
-// DEVICE-LABEL: float @_Z3fooff(float %x, float %y)
-// DEVICE: call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @tex)
-// DEVICE: call %struct.v4f @llvm.nvvm.tex.unified.2d.v4f32.f32(i64 %{{.*}}, float %{{.*}}, float %{{.*}})
-// DEVICE: call i64 @llvm.nvvm.texsurf.handle.internal.p1i64(i64 addrspace(1)* @norm)
-// DEVICE: call %struct.v4f @llvm.nvvm.tex.unified.2d.v4f32.s32(i64 %{{.*}}, i32 %{{.*}}, i32 %{{.*}})
-__attribute__((device)) float foo(float x, float y) {
- return tex2d_ld(tex, x, y).x + tex2d_ld(norm, int(x), int(y)).x;
-}
-
-// HOST: define internal void @[[PREFIX:__cuda]]_register_globals
-// Texture references need registering with correct arguments.
-// HOST: call void @[[PREFIX]]RegisterTexture(i8** %0, i8*{{.*}}({{.*}}@tex{{.*}}), i8*{{.*}}({{.*}}@0{{.*}}), i8*{{.*}}({{.*}}@0{{.*}}), i32 2, i32 0, i32 0)
-// HOST: call void @[[PREFIX]]RegisterTexture(i8** %0, i8*{{.*}}({{.*}}@norm{{.*}}), i8*{{.*}}({{.*}}@1{{.*}}), i8*{{.*}}({{.*}}@1{{.*}}), i32 2, i32 1, i32 0)
-
-// They also need annotating in metadata.
-// DEVICE: !0 = !{i64 addrspace(1)* @tex, !"texture", i32 1}
-// DEVICE: !1 = !{i64 addrspace(1)* @norm, !"texture", i32 1}
diff --git a/clang/test/Misc/pragma-attribute-supported-attributes-list.test b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
index d3705cf9f870..ffef2c717cce 100644
--- a/clang/test/Misc/pragma-attribute-supported-attributes-list.test
+++ b/clang/test/Misc/pragma-attribute-supported-attributes-list.test
@@ -30,8 +30,6 @@
// CHECK-NEXT: CPUSpecific (SubjectMatchRule_function)
// CHECK-NEXT: CUDAConstant (SubjectMatchRule_variable)
// CHECK-NEXT: CUDADevice (SubjectMatchRule_function, SubjectMatchRule_variable)
-// CHECK-NEXT: CUDADeviceBuiltinSurfaceType (SubjectMatchRule_record)
-// CHECK-NEXT: CUDADeviceBuiltinTextureType (SubjectMatchRule_record)
// CHECK-NEXT: CUDAGlobal (SubjectMatchRule_function)
// CHECK-NEXT: CUDAHost (SubjectMatchRule_function)
// CHECK-NEXT: CUDALaunchBounds (SubjectMatchRule_objc_method, SubjectMatchRule_hasType_functionType)
diff --git a/clang/test/SemaCUDA/attr-declspec.cu b/clang/test/SemaCUDA/attr-declspec.cu
index bad86c65c23a..dda12ce8a51f 100644
--- a/clang/test/SemaCUDA/attr-declspec.cu
+++ b/clang/test/SemaCUDA/attr-declspec.cu
@@ -6,19 +6,16 @@
// RUN: %clang_cc1 -DEXPECT_WARNINGS -fms-extensions -fsyntax-only -verify -x c %s
#if defined(EXPECT_WARNINGS)
-// expected-warning at +15 {{'__device__' attribute ignored}}
-// expected-warning at +15 {{'__global__' attribute ignored}}
-// expected-warning at +15 {{'__constant__' attribute ignored}}
-// expected-warning at +15 {{'__shared__' attribute ignored}}
-// expected-warning at +15 {{'__host__' attribute ignored}}
-// expected-warning at +20 {{'__device_builtin_surface_type__' attribute ignored}}
-// expected-warning at +20 {{'__device_builtin_texture_type__' attribute ignored}}
+// expected-warning at +12 {{'__device__' attribute ignored}}
+// expected-warning at +12 {{'__global__' attribute ignored}}
+// expected-warning at +12 {{'__constant__' attribute ignored}}
+// expected-warning at +12 {{'__shared__' attribute ignored}}
+// expected-warning at +12 {{'__host__' attribute ignored}}
//
// (Currently we don't for the other attributes. They are implemented with
// IgnoredAttr, which is ignored irrespective of any LangOpts.)
#else
-// expected-warning at +14 {{'__device_builtin_surface_type__' attribute only applies to classes}}
-// expected-warning at +14 {{'__device_builtin_texture_type__' attribute only applies to classes}}
+// expected-no-diagnostics
#endif
__declspec(__device__) void f_device();
diff --git a/clang/test/SemaCUDA/attributes-on-non-cuda.cu b/clang/test/SemaCUDA/attributes-on-non-cuda.cu
index 215721d273f4..e9e32ce658cb 100644
--- a/clang/test/SemaCUDA/attributes-on-non-cuda.cu
+++ b/clang/test/SemaCUDA/attributes-on-non-cuda.cu
@@ -7,19 +7,16 @@
// RUN: %clang_cc1 -DEXPECT_WARNINGS -fsyntax-only -verify -x c %s
#if defined(EXPECT_WARNINGS)
-// expected-warning at +15 {{'device' attribute ignored}}
-// expected-warning at +15 {{'global' attribute ignored}}
-// expected-warning at +15 {{'constant' attribute ignored}}
-// expected-warning at +15 {{'shared' attribute ignored}}
-// expected-warning at +15 {{'host' attribute ignored}}
-// expected-warning at +21 {{'device_builtin_surface_type' attribute ignored}}
-// expected-warning at +21 {{'device_builtin_texture_type' attribute ignored}}
+// expected-warning at +12 {{'device' attribute ignored}}
+// expected-warning at +12 {{'global' attribute ignored}}
+// expected-warning at +12 {{'constant' attribute ignored}}
+// expected-warning at +12 {{'shared' attribute ignored}}
+// expected-warning at +12 {{'host' attribute ignored}}
//
// NOTE: IgnoredAttr in clang which is used for the rest of
// attributes ignores LangOpts, so there are no warnings.
#else
-// expected-warning at +15 {{'device_builtin_surface_type' attribute only applies to classes}}
-// expected-warning at +15 {{'device_builtin_texture_type' attribute only applies to classes}}
+// expected-no-diagnostics
#endif
__attribute__((device)) void f_device();
diff --git a/clang/test/SemaCUDA/bad-attributes.cu b/clang/test/SemaCUDA/bad-attributes.cu
index d72f74471c23..5aaff9759673 100644
--- a/clang/test/SemaCUDA/bad-attributes.cu
+++ b/clang/test/SemaCUDA/bad-attributes.cu
@@ -70,27 +70,3 @@ void host_fn() {
__device__ void device_fn() {
__constant__ int c; // expected-error {{__constant__ variables must be global}}
}
-
-typedef __attribute__((device_builtin_surface_type)) unsigned long long s0_ty; // expected-warning {{'device_builtin_surface_type' attribute only applies to classes}}
-typedef __attribute__((device_builtin_texture_type)) unsigned long long t0_ty; // expected-warning {{'device_builtin_texture_type' attribute only applies to classes}}
-
-struct __attribute__((device_builtin_surface_type)) s1_ref {}; // expected-error {{illegal device builtin surface reference type 's1_ref' declared here}}
-// expected-note at -1 {{'s1_ref' needs to be instantiated from a class template with proper template arguments}}
-struct __attribute__((device_builtin_texture_type)) t1_ref {}; // expected-error {{illegal device builtin texture reference type 't1_ref' declared here}}
-// expected-note at -1 {{'t1_ref' needs to be instantiated from a class template with proper template arguments}}
-
-template <typename T>
-struct __attribute__((device_builtin_surface_type)) s2_cls_template {}; // expected-error {{illegal device builtin surface reference class template 's2_cls_template' declared here}}
-// expected-note at -1 {{'s2_cls_template' needs to have exactly 2 template parameters}}
-template <typename T>
-struct __attribute__((device_builtin_texture_type)) t2_cls_template {}; // expected-error {{illegal device builtin texture reference class template 't2_cls_template' declared here}}
-// expected-note at -1 {{'t2_cls_template' needs to have exactly 3 template parameters}}
-
-template <int val, void *ptr>
-struct __attribute__((device_builtin_surface_type)) s3_cls_template {}; // expected-error {{illegal device builtin surface reference class template 's3_cls_template' declared here}}
-// expected-note at -1 {{the 1st template parameter of 's3_cls_template' needs to be a type}}
-// expected-note at -2 {{the 2nd template parameter of 's3_cls_template' needs to be an integer or enum value}}
-template <int val, int type, typename T>
-struct __attribute__((device_builtin_texture_type)) t3_cls_template {}; // expected-error {{illegal device builtin texture reference class template 't3_cls_template' declared here}}
-// expected-note at -1 {{the 1st template parameter of 't3_cls_template' needs to be a type}}
-// expected-note at -2 {{the 3rd template parameter of 't3_cls_template' needs to be an integer or enum value}}
diff --git a/llvm/include/llvm/IR/Operator.h b/llvm/include/llvm/IR/Operator.h
index ec594666f0a4..35e08d9215e2 100644
--- a/llvm/include/llvm/IR/Operator.h
+++ b/llvm/include/llvm/IR/Operator.h
@@ -599,25 +599,6 @@ class BitCastOperator
}
};
-class AddrSpaceCastOperator
- : public ConcreteOperator<Operator, Instruction::AddrSpaceCast> {
- friend class AddrSpaceCastInst;
- friend class ConstantExpr;
-
-public:
- Value *getPointerOperand() { return getOperand(0); }
-
- const Value *getPointerOperand() const { return getOperand(0); }
-
- unsigned getSrcAddressSpace() const {
- return getPointerOperand()->getType()->getPointerAddressSpace();
- }
-
- unsigned getDestAddressSpace() const {
- return getType()->getPointerAddressSpace();
- }
-};
-
} // end namespace llvm
#endif // LLVM_IR_OPERATOR_H
More information about the llvm-commits
mailing list