clang/test/SemaHIP/amdgpu-builtin-in-lambda-with-unsupported-attribute.hip

// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -std=c++20 -triple amdgpu6.00 -emit-llvm -fcuda-is-device -verify=no-memrealtime -o - %s
// RUN: %clang_cc1 -std=c++20 -triple amdgpu9.50 -emit-llvm -fcuda-is-device -o - %s

#define __device__ __attribute__((device))
#define __shared__ __attribute__((shared))

struct S {
    static constexpr auto memrealtime_lambda = []() {
        __builtin_amdgcn_s_memrealtime(); // no-memrealtime-error{{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
    };
};

__attribute__((target("s-memrealtime")))
__device__ void test_target_dependant_builtin_attr_fail() {
    S::memrealtime_lambda();
}

constexpr auto memrealtime_lambda = []() {
    __builtin_amdgcn_s_memrealtime(); // no-memrealtime-error{{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
};

__attribute__((target("s-memrealtime")))
__device__ void global_test_target_dependant_builtin_attr_fail() {
    memrealtime_lambda();
}

__attribute__((target("s-memrealtime")))
__device__ void local_test_target_dependant_builtin_attr_fail() {
    static constexpr auto f = []() {
        __builtin_amdgcn_s_memrealtime(); // no-memrealtime-error{{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
    };
    f();
}