1 // Verify the behavior of the +gfxN-insts in the way that
2 // rocm-device-libs should be built with. e.g. If the device libraries has a function
3 // with "+gfx11-insts", that attribute should still be present after linking and not
4 // overwritten with the current target's settings.
5
6 // This is important because at this time, many device-libs functions that are only
7 // available on some GPUs put an attribute such as "+gfx11-insts" so that
8 // AMDGPURemoveIncompatibleFunctions can detect & remove them if needed.
9
10 // Build the fake device library in the way rocm-device-libs should be built.
11 //
12 // RUN: %clang_cc1 -x cl -triple amdgcn-amd-amdhsa\
13 // RUN: -mcode-object-version=none -emit-llvm-bc \
14 // RUN: %S/Inputs/ocml-sample-target-attrs.cl -o %t.bc
15
16 // Check the default behavior
17 // RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 -fcuda-is-device \
18 // RUN: -mlink-builtin-bitcode %t.bc \
19 // RUN: -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,INTERNALIZE
20
21 // RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx1101 -fcuda-is-device \
22 // RUN: -mlink-builtin-bitcode %t.bc -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,INTERNALIZE
23
24 // Check the case where no internalization is performed
25 // RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx803 \
26 // RUN: -fcuda-is-device -mlink-bitcode-file %t.bc -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,NOINTERNALIZE
27
28 // Check the case where no internalization is performed
29 // RUN: %clang_cc1 -x hip -triple amdgcn-amd-amdhsa -target-cpu gfx1101 \
30 // RUN: -fcuda-is-device -mlink-bitcode-file %t.bc -emit-llvm %s -o - | FileCheck %s --check-prefixes=CHECK,NOINTERNALIZE
31
32
33 // CHECK: define {{.*}} i64 @do_intrin_stuff() #[[ATTR:[0-9]+]]
34 // INTERNALIZE: attributes #[[ATTR]] = {{.*}} "target-cpu"="gfx{{.*}}" "target-features"="{{.*}}+gfx11-insts{{.*}}"
35 // NOINTERNALIZE: attributes #[[ATTR]] = {{.*}} "target-features"="+gfx11-insts"
36
37 #define __device__ __attribute__((device))
38 #define __global__ __attribute__((global))
39
40 typedef unsigned long ulong;
41
42 extern "C" {
43 __device__ ulong do_intrin_stuff(void);
44
kernel_f16(ulong * out)45 __global__ void kernel_f16(ulong* out) {
46 *out = do_intrin_stuff();
47 }
48 }
49