yaxunl updated this revision to Diff 204155.
yaxunl marked 5 inline comments as done.
yaxunl added a comment.
Revised by Artem's comments.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D62738/new/
https://reviews.llvm.org/D62738
Files:
include/clang/Basic/Attr.td
include/clang/Basic/AttrDocs.td
lib/AST/TypePrinter.cpp
lib/CodeGen/CodeGenModule.cpp
lib/CodeGen/TargetInfo.cpp
lib/Sema/SemaDeclAttr.cpp
test/AST/ast-dump-cuda-texture.cu
test/CodeGenCUDA/texture.cu
test/SemaCUDA/attr-declspec.cu
test/SemaCUDA/attributes-on-non-cuda.cu
Index: test/SemaCUDA/attributes-on-non-cuda.cu
===================================================================
--- test/SemaCUDA/attributes-on-non-cuda.cu
+++ test/SemaCUDA/attributes-on-non-cuda.cu
@@ -7,11 +7,12 @@
// RUN: %clang_cc1 -DEXPECT_WARNINGS -fsyntax-only -verify -x c %s
#if defined(EXPECT_WARNINGS)
-// expected-warning@+12 {{'device' attribute ignored}}
-// expected-warning@+12 {{'global' attribute ignored}}
-// expected-warning@+12 {{'constant' attribute ignored}}
-// expected-warning@+12 {{'shared' attribute ignored}}
-// expected-warning@+12 {{'host' attribute ignored}}
+// expected-warning@+13 {{'device' attribute ignored}}
+// expected-warning@+13 {{'global' attribute ignored}}
+// expected-warning@+13 {{'constant' attribute ignored}}
+// expected-warning@+13 {{'shared' attribute ignored}}
+// expected-warning@+13 {{'host' attribute ignored}}
+// expected-warning@+20 {{'device_builtin_texture_type' attribute ignored}}
//
// NOTE: IgnoredAttr in clang which is used for the rest of
// attributes ignores LangOpts, so there are no warnings.
Index: test/SemaCUDA/attr-declspec.cu
===================================================================
--- test/SemaCUDA/attr-declspec.cu
+++ test/SemaCUDA/attr-declspec.cu
@@ -6,11 +6,12 @@
// RUN: %clang_cc1 -DEXPECT_WARNINGS -fms-extensions -fsyntax-only -verify -x c %s
#if defined(EXPECT_WARNINGS)
-// expected-warning@+12 {{'__device__' attribute ignored}}
-// expected-warning@+12 {{'__global__' attribute ignored}}
-// expected-warning@+12 {{'__constant__' attribute ignored}}
-// expected-warning@+12 {{'__shared__' attribute ignored}}
-// expected-warning@+12 {{'__host__' attribute ignored}}
+// expected-warning@+13 {{'__device__' attribute ignored}}
+// expected-warning@+13 {{'__global__' attribute ignored}}
+// expected-warning@+13 {{'__constant__' attribute ignored}}
+// expected-warning@+13 {{'__shared__' attribute ignored}}
+// expected-warning@+13 {{'__host__' attribute ignored}}
+// expected-warning@+19 {{'__device_builtin_texture_type__' attribute ignored}}
//
// (Currently we don't for the other attributes. They are implemented with
// IgnoredAttr, which is ignored irrespective of any LangOpts.)
Index: test/CodeGenCUDA/texture.cu
===================================================================
--- /dev/null
+++ test/CodeGenCUDA/texture.cu
@@ -0,0 +1,28 @@
+// REQUIRES: amdgpu-registered-target
+
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -std=c++11 -fcuda-is-device \
+// RUN: -emit-llvm -o - %s | FileCheck -check-prefixes=CUDADEV %s
+// RUN: %clang_cc1 -triple x86_64 -std=c++11 \
+// RUN: -emit-llvm -o - %s | FileCheck -check-prefixes=CUDAHOST %s
+
+// RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -std=c++11 -fvisibility hidden -fapply-global-visibility-to-externs \
+// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefixes=HIPDEV %s
+// RUN: %clang_cc1 -triple x86_64 -std=c++11 \
+// RUN: -emit-llvm -o - -x hip %s | FileCheck -check-prefixes=HIPHOST %s
+
+struct textureReference {
+ int a;
+};
+
+template <class T, int texType, int hipTextureReadMode>
+struct __attribute__((device_builtin_texture_type)) texture : public textureReference {
+texture() { a = 1; }
+};
+
+texture<float, 2, 1> tex;
+// CUDADEV-NOT: @tex
+// CUDAHOST-NOT: call i32 @__hipRegisterVar{{.*}}@tex
+// HIPDEV: @tex = protected addrspace(1) global %struct.texture undef
+// HIPDEV-NOT: declare{{.*}}void @_ZN7textureIfLi2ELi1EEC1Ev
+// HIPHOST: define{{.*}}@_ZN7textureIfLi2ELi1EEC1Ev
+// HIPHOST: call i32 @__hipRegisterVar{{.*}}@tex
Index: test/AST/ast-dump-cuda-texture.cu
===================================================================
--- /dev/null
+++ test/AST/ast-dump-cuda-texture.cu
@@ -0,0 +1,11 @@
+// RUN: %clang_cc1 -fcuda-is-device -ast-dump -ast-dump-filter texture %s | FileCheck -strict-whitespace %s
+// RUN: %clang_cc1 -ast-dump -ast-dump-filter texture %s | FileCheck -strict-whitespace %s
+struct textureReference {
+ int a;
+};
+
+// CHECK: CUDADeviceBuiltinTextureTypeAttr
+template <class T, int texType, int hipTextureReadMode>
+struct __attribute__((device_builtin_texture_type)) texture : public textureReference {
+texture() { a = 1; }
+};
Index: lib/Sema/SemaDeclAttr.cpp
===================================================================
--- lib/Sema/SemaDeclAttr.cpp
+++ lib/Sema/SemaDeclAttr.cpp
@@ -6786,6 +6786,9 @@
case ParsedAttr::AT_CUDAHost:
handleSimpleAttributeWithExclusions<CUDAHostAttr, CUDAGlobalAttr>(S, D, AL);
break;
+ case ParsedAttr::AT_CUDADeviceBuiltinTextureType:
+ handleSimpleAttribute<CUDADeviceBuiltinTextureTypeAttr>(S, D, AL);
+ break;
case ParsedAttr::AT_GNUInline:
handleGNUInlineAttr(S, D, AL);
break;
Index: lib/CodeGen/TargetInfo.cpp
===================================================================
--- lib/CodeGen/TargetInfo.cpp
+++ lib/CodeGen/TargetInfo.cpp
@@ -7848,7 +7848,13 @@
return D->hasAttr<OpenCLKernelAttr>() ||
(isa<FunctionDecl>(D) && D->hasAttr<CUDAGlobalAttr>()) ||
(isa<VarDecl>(D) &&
- (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()));
+ (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
+ (cast<VarDecl>(D)->getType()->isRecordType() &&
+ cast<VarDecl>(D)
+ ->getType()
+ ->getAs<RecordType>()
+ ->getDecl()
+ ->hasAttr<CUDADeviceBuiltinTextureTypeAttr>())));
}
void AMDGPUTargetCodeGenInfo::setTargetAttributes(
Index: lib/CodeGen/CodeGenModule.cpp
===================================================================
--- lib/CodeGen/CodeGenModule.cpp
+++ lib/CodeGen/CodeGenModule.cpp
@@ -2388,6 +2388,12 @@
return ConstantAddress(Aliasee, Alignment);
}
+static bool isCUDATextureType(const QualType &T) {
+ return T->isRecordType() && T->getAs<RecordType>()
+ ->getDecl()
+ ->hasAttr<CUDADeviceBuiltinTextureTypeAttr>();
+}
+
void CodeGenModule::EmitGlobal(GlobalDecl GD) {
const auto *Global = cast<ValueDecl>(GD.getDecl());
@@ -2414,7 +2420,8 @@
if (!Global->hasAttr<CUDADeviceAttr>() &&
!Global->hasAttr<CUDAGlobalAttr>() &&
!Global->hasAttr<CUDAConstantAttr>() &&
- !Global->hasAttr<CUDASharedAttr>())
+ !Global->hasAttr<CUDASharedAttr>() &&
+ !(LangOpts.HIP && isCUDATextureType(Global->getType())))
return;
} else {
// We need to emit host-side 'shadows' for all global
@@ -3769,7 +3776,11 @@
!getLangOpts().CUDAIsDevice &&
(D->hasAttr<CUDAConstantAttr>() || D->hasAttr<CUDADeviceAttr>() ||
D->hasAttr<CUDASharedAttr>());
- if (getLangOpts().CUDA && (IsCUDASharedVar || IsCUDAShadowVar))
+ // HIP texture references have non-trivial default constructors, therefore
+ // they cannot be initialized in device code.
+ if (getLangOpts().CUDA && (IsCUDASharedVar || IsCUDAShadowVar ||
+ (getLangOpts().HIP && getLangOpts().CUDAIsDevice &&
+ isCUDATextureType(D->getType()))))
Init = llvm::UndefValue::get(getTypes().ConvertType(ASTTy));
else if (!InitExpr) {
// This is a tentative definition; tentative definitions are
@@ -3880,7 +3891,8 @@
// global variables become internal definitions. These have to
// be internal in order to prevent name conflicts with global
// host variables with the same name in a different TUs.
- if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>()) {
+ if (D->hasAttr<CUDADeviceAttr>() || D->hasAttr<CUDAConstantAttr>() ||
+ (isCUDATextureType(D->getType()) && LangOpts.HIP)) {
Linkage = llvm::GlobalValue::InternalLinkage;
// Shadow variables and their properties must be registered
Index: lib/AST/TypePrinter.cpp
===================================================================
--- lib/AST/TypePrinter.cpp
+++ lib/AST/TypePrinter.cpp
@@ -1511,6 +1511,7 @@
case attr::SPtr:
case attr::UPtr:
case attr::AddressSpace:
+ case attr::CUDADeviceBuiltinTextureType:
llvm_unreachable("This attribute should have been handled already");
case attr::NSReturnsRetained:
Index: include/clang/Basic/AttrDocs.td
===================================================================
--- include/clang/Basic/AttrDocs.td
+++ include/clang/Basic/AttrDocs.td
@@ -4157,3 +4157,55 @@
``__attribute__((malloc))``.
}];
}
+
+def DeviceBuiltinTextureTypeDocs : Documentation {
+ let Category = DocCatType;
+ let Content = [{
+The GNU style attribute __attribute__((device_builtin_texture_type)) or MSVC
+style attribute __declspec(device_builtin_texture_type) can be added to the
+definition of a class template to indicate it is the HIP texture reference type,
+which is defined as
+
+ .. code-block:: c++
+
+ template <class T, int texType, enum hipTextureReadMode>
+ struct __attribute__((device_builtin_texture_type)) texture
+ : public textureReference { ... }
+
+where T is the data type of texels, texType is the type of the texture, e.g 1D,
+2D, or 3D, hipTextureReadMode is the enumeration type of the read mode of
+the texture. HIP texture reference type is defined as a class template so that
+the template arguments are compile time constants. HIP texture reference also
+contains information which can be changed at run time, e.g. filter mode, address
+mode, which are defined in the base class textureReference. HIP texture
+references can only be defined as global variables. In device code, they can
+be passed to texture fetch API functions to fetch texels. In host code, they
+can be modified and binded to textures by texture binding API functions. They
+should be binded in host code before they can be used in kernels. An example
+of HIP texture reference is as follows:
+
+ .. code-block:: c++
+ texture<float, 2, hipReadModeElementType> tex;
+
+ __global__ void tex2DKernel(float* outputData,
+ int width, int height) {
+ int x = blockIdx.x * blockDim.x + threadIdx.x;
+ int y = blockIdx.y * blockDim.y + threadIdx.y;
+ outputData[y * width + x] = tex2D(tex, x, y);
+ }
+
+ int main() {
+ // ...
+ tex.addressMode[0] = hipAddressModeWrap;
+ tex.addressMode[1] = hipAddressModeWrap;
+ tex.filterMode = hipFilterModePoint;
+ tex.normalized = 0;
+ hipBindTextureToArray(tex, hipArray, channelDesc);
+ tex2DKernel<<<dim3(dimGrid), dim3(dimBlock), 0, 0>>>(dData, width, height);
+ // ...
+ }
+}
+
+It is ignored for CUDA and other languages.
+ }];
+}
\ No newline at end of file
Index: include/clang/Basic/Attr.td
===================================================================
--- include/clang/Basic/Attr.td
+++ include/clang/Basic/Attr.td
@@ -962,10 +962,11 @@
let LangOpts = [CUDA];
}
-def CUDADeviceBuiltinTextureType : IgnoredAttr {
+def CUDADeviceBuiltinTextureType : TypeAttr {
let Spellings = [GNU<"device_builtin_texture_type">,
Declspec<"__device_builtin_texture_type__">];
let LangOpts = [CUDA];
+ let Documentation = [DeviceBuiltinTextureTypeDocs];
}
def CUDAGlobal : InheritableAttr {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
