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


1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34

// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -std=c++20 -triple amdgcn -target-cpu tahiti -emit-llvm -fcuda-is-device -verify=no-memrealtime -o - %s
// RUN: %clang_cc1 -std=c++20 -triple amdgcn -target-cpu gfx950 -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();
}