c7b683c126
- Skip generating profile data on `__global__` function in the host
compilation. It's a host-side stub function only and don't have
profile instrumentation generated on the real function body. The extra
profile data results in the malformed instrumentation profile data.
- Skip generating region mapping on functions in the wrong-side, i.e.,
+ For the device compilation, skip host-only functions; and,
+ For the host compilation, skip device-only functions (including
`__global__` functions.)
- As the device-side profiling is not ready yet, only host-side profile
code generation is checked.
Differential Revision: https://reviews.llvm.org/D85276
21 lines
994 B
Plaintext
21 lines
994 B
Plaintext
// RUN: echo "GPU binary would be here" > %t
|
|
// RUN: %clang_cc1 -fprofile-instrument=clang -triple x86_64-linux-gnu -target-sdk-version=8.0 -fcuda-include-gpubinary %t -emit-llvm -o - %s | FileCheck --check-prefix=PGOGEN %s
|
|
// RUN: %clang_cc1 -fprofile-instrument=clang -fcoverage-mapping -triple x86_64-linux-gnu -target-sdk-version=8.0 -fcuda-include-gpubinary %t -emit-llvm -o - %s | FileCheck --check-prefix=COVMAP %s
|
|
// RUN: %clang_cc1 -fprofile-instrument=clang -fcoverage-mapping -dump-coverage-mapping -triple x86_64-linux-gnu -target-sdk-version=8.0 -fcuda-include-gpubinary %t -emit-llvm-only -o - %s | FileCheck --check-prefix=MAPPING %s
|
|
|
|
#include "Inputs/cuda.h"
|
|
|
|
// PGOGEN-NOT: @__profn_{{.*kernel.*}} =
|
|
// COVMAP-COUNT-2: section "__llvm_covfun", comdat
|
|
// COVMAP-NOT: section "__llvm_covfun", comdat
|
|
// MAPPING-NOT: {{.*dfn.*}}:
|
|
// MAPPING-NOT: {{.*kernel.*}}:
|
|
|
|
__device__ void dfn(int i) {}
|
|
|
|
__global__ void kernel(int i) { dfn(i); }
|
|
|
|
void host(void) {
|
|
kernel<<<1, 1>>>(1);
|
|
}
|