Summary:
- Use `device_builtin_surface` and `device_builtin_texture` for
surface/texture reference support. So far, both the host and device
use the same reference type, which could be revised later when
interface/implementation is stablized.
Reviewers: yaxunl
Subscribers: cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D77583
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
This makes it possible for plugin attributes to actually do something, and also
removes a lot of boilerplate for simple attributes in SemaDeclAttr.cpp.
Differential Revision: https://reviews.llvm.org/D31342
The SVE ACLE allows using a short-form for the intrinsics, e.g.
the following two declarations generate the same code:
svuint32_t svld1(svbool_t, uint32_t const *);
svuint32_t svld1_u32(svbool_t, uint32_t const *);
using the attribute:
__clang_arm_builtin_alias
so that any call to svld1(svbool_t, uint32_t const *) will
map to __builtin_sve_svld1_u32.
Reviewers: SjoerdMeijer, miyuki, efriedma, simon_tatham, rengolin
Reviewed By: SjoerdMeijer
Tags: #clang
Differential Revision: https://reviews.llvm.org/D75861
Summary:
[Clang] Attribute to allow defining undef global variables
Initializing global variables is very cheap on hosted implementations. The
C semantics of zero initializing globals work very well there. It is not
necessarily cheap on freestanding implementations. Where there is no loader
available, code must be emitted near the start point to write the appropriate
values into memory.
At present, external variables can be declared in C++ and definitions provided
in assembly (or IR) to achive this effect. This patch provides an attribute in
order to remove this reason for writing assembly for performance sensitive
freestanding implementations.
A close analogue in tree is LDS memory for amdgcn, where the kernel is
responsible for initializing the memory after it starts executing on the gpu.
Uninitalized variables in LDS are observably cheaper than zero initialized.
Patch is loosely based on the cuda __shared__ and opencl __local variable
implementation which also produces undef global variables.
Reviewers: kcc, rjmccall, rsmith, glider, vitalybuka, pcc, eugenis, vlad.tsyrklevich, jdoerfert, gregrodgers, jfb, aaron.ballman
Reviewed By: rjmccall, aaron.ballman
Subscribers: Anastasia, aaron.ballman, davidb, Quuxplusone, dexonsmith, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D74361
A previous patch rejected alignof for sizeless types. This patch
extends that to cover the "aligned" attribute and _Alignas. Since
sizeless types are not meant to be used for long-term data, cannot
be used in aggregates, and cannot have static storage duration,
there shouldn't be any need to fiddle with their alignment.
Like with alignof, this is a conservative position that can be
relaxed in future if it turns out to be too restrictive.
Differential Revision: https://reviews.llvm.org/D75573
Summary:
This patch generalizes the existing code to support CDE intrinsics
which will share some properties with existing MVE intrinsics
(some of the intrinsics will be polymorphic and accept/return values
of MVE vector types).
Specifically the patch:
* Adds new tablegen backends -gen-arm-cde-builtin-def,
-gen-arm-cde-builtin-codegen, -gen-arm-cde-builtin-sema,
-gen-arm-cde-builtin-aliases, -gen-arm-cde-builtin-header based on
existing MVE backends.
* Renames the '__clang_arm_mve_alias' attribute into
'__clang_arm_builtin_alias' (it will be used with CDE intrinsics as
well as MVE intrinsics)
* Implements semantic checks for the coprocessor argument of the CDE
intrinsics as well as the existing coprocessor intrinsics.
* Adds one CDE intrinsic __arm_cx1 to test the above changes
Reviewers: simon_tatham, MarkMurrayARM, ostannard, dmgreen
Reviewed By: simon_tatham
Subscribers: sdesmalen, mgorny, kristof.beyls, danielkiss, cfe-commits, llvm-commits
Tags: #clang, #llvm
Differential Revision: https://reviews.llvm.org/D75850
Reduces compile time of SemaDeclAttr.cpp down to 28s from 50s. The new
TU does a few RecursiveASTVisitor instantiations, so it takes 30s.
Reviewed By: rsmith
Differential Revision: https://reviews.llvm.org/D73385
There is llvm::Value::MaximumAlignment, which is numerically
equivalent to these constants, but we can't use it directly
because we can't include llvm IR headers in clang Sema.
So instead, copy-paste the constant, and fixup the places to use it.
This was initially reviewed in https://reviews.llvm.org/D72998
Summary:
For `__builtin_assume_aligned()`, we do validate that the alignment
is not greater than `536870912` (D68824), but we don't do that for
`__attribute__((assume_aligned(N)))` attribute.
I suspect we should.
This was initially committed in a4cfb15d15
but reverted in 210f0882c9 due to
suspicious bot failures.
Reviewers: erichkeane, aaron.ballman, hfinkel, rsmith, jdoerfert
Reviewed By: erichkeane
Subscribers: cfe-commits, llvm-commits
Tags: #llvm, #clang
Differential Revision: https://reviews.llvm.org/D72994
Summary:
I kind-of understand why it is restricted to integer-typed arguments,
for general enum's the value passed is not nessesairly the alignment implied,
although one might say that user would know best.
But we clearly should whitelist `std::align_val_t`,
which is just a thin wrapper over `std::size_t`,
and is the C++ standard way of specifying alignment.
Reviewers: erichkeane, rsmith, aaron.ballman, jdoerfert
Reviewed By: erichkeane
Subscribers: cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D73019
Summary:
I initially encountered those assertions when trying to create
this IR `alignment` attribute from clang's `__attribute__((assume_aligned(imm)))`,
because until D72994 there is no sanity checking for the value of `imm`.
But even then, we have `llvm::Value::MaximumAlignment` constant (which is `536870912`),
which is enforced for clang attributes, and then there are some other magical constant
(`0x40000000` i.e. `1073741824` i.e. `2 * 536870912`) in
`Attribute::getWithAlignment()`/`AttrBuilder::addAlignmentAttr()`.
I strongly suspect that `0x40000000` is incorrect,
and that also should be `llvm::Value::MaximumAlignment`.
Reviewers: erichkeane, hfinkel, jdoerfert, gchatelet, courbet
Reviewed By: erichkeane
Subscribers: hiraditya, cfe-commits, llvm-commits
Tags: #llvm, #clang
Differential Revision: https://reviews.llvm.org/D72998
Summary:
For `__builtin_assume_aligned()`, we do validate that the alignment
is not greater than `536870912` (D68824), but we don't do that for
`__attribute__((assume_aligned(N)))` attribute.
I suspect we should.
Reviewers: erichkeane, aaron.ballman, hfinkel, rsmith, jdoerfert
Reviewed By: erichkeane
Subscribers: cfe-commits, llvm-commits
Tags: #llvm, #clang
Differential Revision: https://reviews.llvm.org/D72994
This feature is generic. Make it applicable for AArch64 and X86 because
the backend has only implemented NOP insertion for AArch64 and X86.
Reviewed By: nickdesaulniers, aaron.ballman
Differential Revision: https://reviews.llvm.org/D72221
Summary:
Avoid using the `nocf_check` attribute with Control Flow Guard. Instead, use a
new `"guard_nocf"` function attribute to indicate that checks should not be
added on indirect calls within that function. Add support for
`__declspec(guard(nocf))` following the same syntax as MSVC.
Reviewers: rnk, dmajor, pcc, hans, aaron.ballman
Reviewed By: aaron.ballman
Subscribers: aaron.ballman, tomrittervg, hiraditya, cfe-commits, llvm-commits
Tags: #clang, #llvm
Differential Revision: https://reviews.llvm.org/D72167
Function trailing requires clauses now parsed, supported in overload resolution and when calling, referencing and taking the address of functions or function templates.
Differential Revision: https://reviews.llvm.org/D43357
It turns out it is useful to be able to define the deref type as void.
In case we have a type erased owner, we want to express that the pointee
can be basically any type. It should not be unnatural to have a void
deref type as we already familiar with "pointers to void".
Differential Revision: https://reviews.llvm.org/D72097
msvc allows a subsequent declaration of a uuid attribute on a
struct/class. Mirror this behavior in clang-cl.
Reviewed By: rnk
Differential Revision: https://reviews.llvm.org/D71439
These annotations will be used in an upcomming static analyzer check
that finds handle leaks, use after releases, and double releases.
Differential Revision: https://reviews.llvm.org/D70469
This change updates the clang front end to add symbols to llvm.used
when they have explicit export_name attribute.
Differential Revision: https://reviews.llvm.org/D71493
This is equivalent to the existing `import_name` and `import_module`
attributes which control the import names in the final wasm binary
produced by lld.
This maps the existing
This attribute currently requires a string rather than using the
symbol name for a couple of reasons:
1. Avoid confusion with static and dynamic linking which is
based on symbol name. Exporting a function from a wasm module using
this directive is orthogonal to both static and dynamic linking.
2. Avoids name mangling.
Differential Revision: https://reviews.llvm.org/D70520
SYCL is single source offload programming model relying on compiler to
separate device code (i.e. offloaded to an accelerator) from the code
executed on the host.
Here is code example of the SYCL program to demonstrate compiler
outlining work:
```
int foo(int x) { return ++x; }
int bar(int x) { throw std::exception("CPU code only!"); }
...
using namespace cl::sycl;
queue Q;
buffer<int, 1> a(range<1>{1024});
Q.submit([&](handler& cgh) {
auto A = a.get_access<access::mode::write>(cgh);
cgh.parallel_for<init_a>(range<1>{1024}, [=](id<1> index) {
A[index] = index[0] + foo(42);
});
}
...
```
SYCL device compiler must compile lambda expression passed to
cl::sycl::handler::parallel_for method and function foo called from this
lambda expression for an "accelerator". SYCL device compiler also must
ignore bar function as it's not required for offloaded code execution.
This patch adds the sycl_kernel attribute, which is used to mark code
passed to cl::sycl::handler::parallel_for as "accelerated code".
Attribute must be applied to function templates which parameters include
at least "kernel name" and "kernel function object". These parameters
will be used to establish an ABI between the host application and
offloaded part.
Reviewers: jlebar, keryell, Naghasan, ABataev, Anastasia, bader, aaron.ballman, rjmccall, rsmith
Reviewed By: keryell, bader
Subscribers: mgorny, OlegM, ArturGainullin, agozillon, aaron.ballman, ebevhan, Anastasia, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D60455
Signed-off-by: Alexey Bader <alexey.bader@intel.com>
This avoids the need to include Attr.h in DeclCXX.h for a four-value
enum. Removing the include will be done separately, since it is large
and risky change.
__attribute__((objc_direct)) is an attribute on methods declaration, and
__attribute__((objc_direct_members)) on implementation, categories or
extensions.
A `direct` property specifier is added (@property(direct) type name)
These attributes / specifiers cause the method to have no associated
Objective-C metadata (for the property or the method itself), and the
calling convention to be a direct C function call.
The symbol for the method has enforced hidden visibility and such direct
calls are hence unreachable cross image. An explicit C function must be
made if so desired to wrap them.
The implicit `self` and `_cmd` arguments are preserved, however to
maintain compatibility with the usual `objc_msgSend` semantics,
3 fundamental precautions are taken:
1) for instance methods, `self` is nil-checked. On arm64 backends this
typically adds a single instruction (cbz x0, <closest-ret>) to the
codegen, for the vast majority of the cases when the return type is a
scalar.
2) for class methods, because the class may not be realized/initialized
yet, a call to `[self self]` is emitted. When the proper deployment
target is used, this is optimized to `objc_opt_self(self)`.
However, long term we might want to emit something better that the
optimizer can reason about. When inlining kicks in, these calls
aren't optimized away as the optimizer has no idea that a single call
is really necessary.
3) the calling convention for the `_cmd` argument is changed: the caller
leaves the second argument to the call undefined, and the selector is
loaded inside the body when it's referenced only.
As far as error reporting goes, the compiler refuses:
- making any overloads direct,
- making an overload of a direct method,
- implementations marked as direct when the declaration in the
interface isn't (the other way around is allowed, as the direct
attribute is inherited from the declaration),
- marking methods required for protocol conformance as direct,
- messaging an unqualified `id` with a direct method,
- forming any @selector() expression with only direct selectors.
As warnings:
- any inconsistency of direct-related calling convention when
@selector() or messaging is used,
- forming any @selector() expression with a possibly direct selector.
Lastly an `objc_direct_members` attribute is added that can decorate
`@implementation` blocks and causes methods only declared there (and in
no `@interface`) to be automatically direct. When decorating an
`@interface` then all methods and properties declared in this block are
marked direct.
Radar-ID: rdar://problem/2684889
Differential Revision: https://reviews.llvm.org/D69991
Reviewed-By: John McCall
This is a resubmission for the previous reverted commit
9434360401 with the same subject. This commit fixed the
segfault issue and addressed additional review comments.
This patch introduced a new bpf specific attribute which can
be added to struct or union definition. For example,
struct s { ... } __attribute__((preserve_access_index));
union u { ... } __attribute__((preserve_access_index));
The goal is to simplify user codes for cases
where preserve access index happens for certain struct/union,
so user does not need to use clang __builtin_preserve_access_index
for every members.
The attribute has no effect if -g is not specified.
When the attribute is specified and -g is specified, any member
access defined by that structure or union, including array subscript
access and inner records, will be preserved through
__builtin_preserve_{array,struct,union}_access_index()
IR intrinsics, which will enable relocation generation
in bpf backend.
The following is an example to illustrate the usage:
-bash-4.4$ cat t.c
#define __reloc__ __attribute__((preserve_access_index))
struct s1 {
int c;
} __reloc__;
struct s2 {
union {
struct s1 b[3];
};
} __reloc__;
struct s3 {
struct s2 a;
} __reloc__;
int test(struct s3 *arg) {
return arg->a.b[2].c;
}
-bash-4.4$ clang -target bpf -g -S -O2 t.c
A relocation with access string "0:0:0:0:2:0" will be generated
representing access offset of arg->a.b[2].c.
forward declaration with attribute is also handled properly such
that the attribute is copied and populated in real record definition.
Differential Revision: https://reviews.llvm.org/D69759
This patch introduced a new bpf specific attribute which can
be added to struct or union definition. For example,
struct s { ... } __attribute__((preserve_access_index));
union u { ... } __attribute__((preserve_access_index));
The goal is to simplify user codes for cases
where preserve access index happens for certain struct/union,
so user does not need to use clang __builtin_preserve_access_index
for every members.
The attribute has no effect if -g is not specified.
When the attribute is specified and -g is specified, any member
access defined by that structure or union, including array subscript
access and inner records, will be preserved through
__builtin_preserve_{array,struct,union}_access_index()
IR intrinsics, which will enable relocation generation
in bpf backend.
The following is an example to illustrate the usage:
-bash-4.4$ cat t.c
#define __reloc__ __attribute__((preserve_access_index))
struct s1 {
int c;
} __reloc__;
struct s2 {
union {
struct s1 b[3];
};
} __reloc__;
struct s3 {
struct s2 a;
} __reloc__;
int test(struct s3 *arg) {
return arg->a.b[2].c;
}
-bash-4.4$ clang -target bpf -g -S -O2 t.c
A relocation with access string "0:0:0:0:2:0" will be generated
representing access offset of arg->a.b[2].c.
forward declaration with attribute is also handled properly such
that the attribute is copied and populated in real record definition.
Differential Revision: https://reviews.llvm.org/D69759
Summary:
This is a follow up on https://reviews.llvm.org/D61634
This patch is simpler and only adds the no_builtin attribute.
Reviewers: tejohnson, courbet, theraven, t.p.northover, jdoerfert
Subscribers: mgrang, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D68028
This is a re-submit after it got reverted in https://reviews.llvm.org/rGbd8791610948 since the breakage doesn't seem to come from this patch.
Summary:
This is a follow up on https://reviews.llvm.org/D61634
This patch is simpler and only adds the no_builtin attribute.
Reviewers: tejohnson, courbet, theraven, t.p.northover, jdoerfert
Subscribers: mgrang, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D68028
This commit sets up the infrastructure for auto-generating <arm_mve.h>
and doing clang-side code generation for the builtins it relies on,
and demonstrates that it works by implementing a representative sample
of the ACLE intrinsics, more or less matching the ones introduced in
LLVM IR by D67158,D68699,D68700.
Like NEON, that header file will provide a set of vector types like
uint16x8_t and C functions with names like vaddq_u32(). Unlike NEON,
the ACLE spec for <arm_mve.h> includes a polymorphism system, so that
you can write plain vaddq() and disambiguate by the vector types you
pass to it.
Unlike the corresponding NEON code, I've arranged to make every user-
facing ACLE intrinsic into a clang builtin, and implement all the code
generation inside clang. So <arm_mve.h> itself contains nothing but
typedefs and function declarations, with the latter all using the new
`__attribute__((__clang_builtin))` system to arrange that the user-
facing function names correspond to the right internal BuiltinIDs.
So the new MveEmitter tablegen system specifies the full sequence of
IRBuilder operations that each user-facing ACLE intrinsic should
translate into. Where possible, the ACLE intrinsics map to standard IR
operations such as vector-typed `add` and `fadd`; where no standard
representation exists, I call down to the sample IR intrinsics
introduced in an earlier commit.
Doing it like this means that you get the polymorphism for free just
by using __attribute__((overloadable)): the clang overload resolution
decides which function declaration is the relevant one, and _then_ its
BuiltinID is looked up, so by the time we're doing code generation,
that's all been resolved by the standard system. It also means that
you get really nice error messages if the user passes the wrong
combination of types: clang will show the declarations from the header
file and explain why each one doesn't match.
(The obvious alternative approach would be to have wrapper functions
in <arm_mve.h> which pass their arguments to the underlying builtins.
But that doesn't work in the case where one of the arguments has to be
a constant integer: the wrapper function can't pass the constantness
through. So you'd have to do that case using a macro instead, and then
use C11 `_Generic` to handle the polymorphism. Then you have to add
horrible workarounds because `_Generic` requires even the untaken
branches to type-check successfully, and //then// if the user gets the
types wrong, the error message is totally unreadable!)
Reviewers: dmgreen, miyuki, ostannard
Subscribers: mgorny, javed.absar, kristof.beyls, cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D67161
This allows you to declare a function with a name of your choice (say
`foo`), but have clang treat it as if it were a builtin function (say
`__builtin_foo`), by writing
static __inline__ __attribute__((__clang_arm_mve_alias(__builtin_foo)))
int foo(args);
I'm intending to use this for the ACLE intrinsics for MVE, which have
to be polymorphic on their argument types and also need to be
implemented by builtins. To avoid having to implement the polymorphism
with several layers of nested _Generic and make error reporting
hideous, I want to make all the user-facing intrinsics correspond
directly to clang builtins, so that after clang resolves
__attribute__((overloadable)) polymorphism it's already holding the
right BuiltinID for the intrinsic it selected.
However, this commit itself just introduces the new attribute, and
doesn't use it for anything.
To avoid unanticipated side effects if this attribute is used to make
aliases to other builtins, there's a restriction mechanism: only
(BuiltinID, alias) pairs that are approved by the function
ArmMveAliasValid() will be permitted. At present, that function
doesn't permit anything, because the Tablegen that will generate its
list of valid pairs isn't yet implemented. So the only test of this
facility is one that checks that an unapproved builtin _can't_ be
aliased.
Reviewers: dmgreen, miyuki, ostannard
Subscribers: cfe-commits
Tags: #clang
Differential Revision: https://reviews.llvm.org/D67159
The static analyzer is warning about potential null dereferences, but in these cases we should be able to use castAs<> directly and if not assert will fire for us.
llvm-svn: 373753
The static analyzer is warning about potential null dereferences, but in these cases we should be able to use castAs<RecordType> directly and if not assert will fire for us.
llvm-svn: 373584
This matches how GCC handles it, see e.g. https://gcc.godbolt.org/z/HPplnl.
GCC documents the gnu_inline attribute with "In C++, this attribute does
not depend on extern in any way, but it still requires the inline keyword
to enable its special behavior."
The previous behaviour of gnu_inline in C++, without the extern
keyword, can be traced back to the original commit that added
support for gnu_inline, SVN r69045.
Differential Revision: https://reviews.llvm.org/D67414
llvm-svn: 373078
Summary:
- Even though only `void` is still accepted as the deduced return type,
enabling deduction/instantiation on the return type allows more
consistent coding.
Reviewers: tra, jlebar
Subscribers: cfe-commits, yaxunl
Tags: #clang
Differential Revision: https://reviews.llvm.org/D68031
llvm-svn: 372898
Michael Liao