26 lines · cpp
1// REQUIRES: amdgpu-registered-target2// Test that the accelerator code selection pass only gets invoked after linking3 4// Ensure Pass HipStdParAcceleratorCodeSelectionPass is not invoked in PreLink.5// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -mllvm -amdgpu-enable-hipstdpar -flto -emit-llvm-bc -fcuda-is-device -fdebug-pass-manager \6// RUN: %s -o /dev/null 2>&1 | FileCheck --check-prefix=HIPSTDPAR-PRE %s7// HIPSTDPAR-PRE: Running pass: EntryExitInstrumenterPass8// HIPSTDPAR-PRE-NEXT: Running pass: EntryExitInstrumenterPass9// HIPSTDPAR-PRE-NOT: Running pass: HipStdParAcceleratorCodeSelectionPass10// HIPSTDPAR-PRE-NEXT: Running pass: AlwaysInlinerPass11 12// Ensure Pass HipStdParAcceleratorCodeSelectionPass is invoked in PostLink.13// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -mllvm -amdgpu-enable-hipstdpar -fcuda-is-device -fdebug-pass-manager -emit-llvm \14// RUN: %s -o /dev/null 2>&1 | FileCheck --check-prefix=HIPSTDPAR-POST %s15// HIPSTDPAR-POST: Running pass: HipStdParAcceleratorCodeSelection16 17#define __device__ __attribute__((device))18 19void foo(float *a, float b) {20 *a = b;21}22 23__device__ void bar(float *a, float b) {24 *a = b;25}26