File tree Expand file tree Collapse file tree 1 file changed +29
-0
lines changed
Expand file tree Collapse file tree 1 file changed +29
-0
lines changed Original file line number Diff line number Diff line change @@ -5033,6 +5033,35 @@ target specific type that behaves as if its C++ definition was the following:
50335033 operator bool () const noexcept;
50345034 };
50355035
5036+ The builtins can be used in C as well, wherein the
5037+ ``__amdgpu_feature_predicate_t `` type behaves as an opaque, forward declared
5038+ type with conditional automated conversion to ``_Bool `` when used as the
5039+ predicate argument to a control structure:
5040+
5041+ .. code-block :: c
5042+
5043+ struct __amdgpu_feature_predicate_t ret(); // Error
5044+ void arg(struct __amdgpu_feature_predicate_t); // Error
5045+ void local() {
5046+ struct __amdgpu_feature_predicate_t x; // Error
5047+ struct __amdgpu_feature_predicate_t y =
5048+ __builtin_amdgcn_processor_is("gfx900"); // Error
5049+ }
5050+ void valid_use() {
5051+ _Bool x = (_Bool)__builtin_amdgcn_processor_is("gfx900"); // OK
5052+ if (__builtin_amdgcn_processor_is("gfx900")) // Implicit cast to _Bool
5053+ return;
5054+ for (; __builtin_amdgcn_processor_is("gfx900");) // Implicit cast to _Bool
5055+ break;
5056+ while (__builtin_amdgcn_processor_is("gfx900")) // Implicit cast to _Bool
5057+ break;
5058+ do {
5059+ break;
5060+ } while (__builtin_amdgcn_processor_is("gfx900")); // Implicit cast to _Bool
5061+
5062+ __builtin_amdgcn_processor_is("gfx900") ? x : !x;
5063+ }
5064+
50365065 The boolean interpretation of the predicate values returned by the builtins:
50375066
50385067* indicates whether the current target matches the argument; the argument MUST
You can’t perform that action at this time.
0 commit comments