35 lines · plain
1// REQUIRES: amdgpu-registered-target2// RUN: %clang_cc1 -std=c++20 -triple amdgcn -target-cpu tahiti -emit-llvm -fcuda-is-device -verify=no-memrealtime -o - %s3// RUN: %clang_cc1 -std=c++20 -triple amdgcn -target-cpu gfx950 -emit-llvm -fcuda-is-device -o - %s4 5#define __device__ __attribute__((device))6#define __shared__ __attribute__((shared))7 8struct S {9 static constexpr auto memrealtime_lambda = []() {10 __builtin_amdgcn_s_memrealtime(); // no-memrealtime-error{{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}11 };12};13 14__attribute__((target("s-memrealtime")))15__device__ void test_target_dependant_builtin_attr_fail() {16 S::memrealtime_lambda();17}18 19constexpr auto memrealtime_lambda = []() {20 __builtin_amdgcn_s_memrealtime(); // no-memrealtime-error{{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}21};22 23__attribute__((target("s-memrealtime")))24__device__ void global_test_target_dependant_builtin_attr_fail() {25 memrealtime_lambda();26}27 28__attribute__((target("s-memrealtime")))29__device__ void local_test_target_dependant_builtin_attr_fail() {30 static constexpr auto f = []() {31 __builtin_amdgcn_s_memrealtime(); // no-memrealtime-error{{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}32 };33 f();34}35