49 lines · plain
1// Verify the behavior of the +gfxN-insts in the way that2// rocm-device-libs should be built with. e.g. If the device libraries has a function3// with "+gfx11-insts", that attribute should still be present after linking and not4// overwritten with the current target's settings.5 6// This is important because at this time, many device-libs functions that are only7// available on some GPUs put an attribute such as "+gfx11-insts" so that8// 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.bc15 16// Check the default behavior17// 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,INTERNALIZE20 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,INTERNALIZE23 24// Check the case where no internalization is performed25// 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,NOINTERNALIZE27 28// Check the case where no internalization is performed29// 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,NOINTERNALIZE31 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 40typedef unsigned long ulong;41 42extern "C" {43__device__ ulong do_intrin_stuff(void);44 45__global__ void kernel_f16(ulong* out) {46 *out = do_intrin_stuff();47 }48}49