clang/test/CodeGenHipStdPar/select-accelerator-code-pass-ordering.cpp


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

// REQUIRES: amdgpu-registered-target
// Test that the accelerator code selection pass only gets invoked after linking

// Ensure Pass HipStdParAcceleratorCodeSelectionPass is not invoked in PreLink.
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -mllvm -amdgpu-enable-hipstdpar -flto -emit-llvm-bc -fcuda-is-device -fdebug-pass-manager \
// RUN:  %s -o /dev/null 2>&1 | FileCheck --check-prefix=HIPSTDPAR-PRE %s
// HIPSTDPAR-PRE: Running pass: EntryExitInstrumenterPass
// HIPSTDPAR-PRE-NEXT: Running pass: EntryExitInstrumenterPass
// HIPSTDPAR-PRE-NOT: Running pass: HipStdParAcceleratorCodeSelectionPass
// HIPSTDPAR-PRE-NEXT: Running pass: AlwaysInlinerPass

// Ensure Pass HipStdParAcceleratorCodeSelectionPass is invoked in PostLink.
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -mllvm -amdgpu-enable-hipstdpar -fcuda-is-device -fdebug-pass-manager -emit-llvm \
// RUN:  %s -o /dev/null 2>&1 | FileCheck --check-prefix=HIPSTDPAR-POST %s
// HIPSTDPAR-POST: Running pass: HipStdParAcceleratorCodeSelection

#define __device__ __attribute__((device))

void foo(float *a, float b) {
  *a = b;
}

__device__ void bar(float *a, float b) {
  *a = b;
}