maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity
llvm-project/clang/test/CodeGenCUDA/ptx-kernels.cu

Ignoring revisions in .git-blame-ignore-revs. Click here to bypass and see the normal blame view.

36 lines
1.1 KiB
Plaintext
Raw Normal View History

[CUDA] Add implicit __attribute__((used)) to all __global__ functions. This makes sure that we emit kernels that were instantiated from the host code and which would never be explicitly referenced by anything else on device side. Differential Revision: http://reviews.llvm.org/D11666 llvm-svn: 248293
2015-09-23 01:22:51 +08:00
// Make sure that __global__ functions are emitted along with correct
// annotations and are added to @llvm.used to prevent their elimination.
// REQUIRES: nvptx-registered-target
//
Replace PTX back-end with NVPTX back-end in all places where Clang cares NV_CONTRIB llvm-svn: 157403
2012-05-25 01:43:12 +08:00
// RUN: %clang_cc1 %s -triple nvptx-unknown-unknown -fcuda-is-device -emit-llvm -o - | FileCheck %s
CUDA: set proper calling conventions for PTX llvm-svn: 141296
2011-10-07 00:49:54 +08:00
Move all CUDA testing inputs to Inputs/ subdirectory inside the tests. llvm-svn: 207453
2014-04-29 06:21:28 +08:00
#include "Inputs/cuda.h"
CUDA: set proper calling conventions for PTX llvm-svn: 141296
2011-10-07 00:49:54 +08:00
[test] Add {{.*}} to make ELF tests immune to dso_local/dso_preemptable/(none) differences For a default visibility external linkage definition, dso_local is set for ELF -fno-pic/-fpie and COFF and Mach-O. Since default clang -cc1 for ELF is similar to -fpic ("PIC Level" is not set), this nuance causes unneeded binary format differences. To make emitted IR similar, ELF -cc1 -fpic will default to -fno-semantic-interposition, which sets dso_local for default visibility external linkage definitions. To make this flip smooth and enable future (dso_local as definition default), this patch replaces (function) `define ` with `define{{.*}} `, (variable/constant/alias) `= ` with `={{.*}} `, or inserts appropriate `{{.*}} `.
2020-12-31 16:27:11 +08:00
// CHECK-LABEL: define{{.*}} void @device_function
Use kernel metadata to differentiate between kernel and device functions for the NVPTX target. llvm-svn: 178418
2013-03-30 22:38:24 +08:00
extern "C"
CUDA: set proper calling conventions for PTX llvm-svn: 141296
2011-10-07 00:49:54 +08:00
__device__ void device_function() {}
[test] Add {{.*}} to make ELF tests immune to dso_local/dso_preemptable/(none) differences For a default visibility external linkage definition, dso_local is set for ELF -fno-pic/-fpie and COFF and Mach-O. Since default clang -cc1 for ELF is similar to -fpic ("PIC Level" is not set), this nuance causes unneeded binary format differences. To make emitted IR similar, ELF -cc1 -fpic will default to -fno-semantic-interposition, which sets dso_local for default visibility external linkage definitions. To make this flip smooth and enable future (dso_local as definition default), this patch replaces (function) `define ` with `define{{.*}} `, (variable/constant/alias) `= ` with `={{.*}} `, or inserts appropriate `{{.*}} `.
2020-12-31 16:27:11 +08:00
// CHECK-LABEL: define{{.*}} void @global_function
Use kernel metadata to differentiate between kernel and device functions for the NVPTX target. llvm-svn: 178418
2013-03-30 22:38:24 +08:00
extern "C"
CUDA: set proper calling conventions for PTX llvm-svn: 141296
2011-10-07 00:49:54 +08:00
__global__ void global_function() {
Use kernel metadata to differentiate between kernel and device functions for the NVPTX target. llvm-svn: 178418
2013-03-30 22:38:24 +08:00
// CHECK: call void @device_function
CUDA: set proper calling conventions for PTX llvm-svn: 141296
2011-10-07 00:49:54 +08:00
device_function();
}
Use kernel metadata to differentiate between kernel and device functions for the NVPTX target. llvm-svn: 178418
2013-03-30 22:38:24 +08:00
[CUDA] Add implicit __attribute__((used)) to all __global__ functions. This makes sure that we emit kernels that were instantiated from the host code and which would never be explicitly referenced by anything else on device side. Differential Revision: http://reviews.llvm.org/D11666 llvm-svn: 248293
2015-09-23 01:22:51 +08:00
// Make sure host-instantiated kernels are preserved on device side.
template <typename T> __global__ void templated_kernel(T param) {}
[test] Add {{.*}} to make ELF tests immune to dso_local/dso_preemptable/(none) differences For a default visibility external linkage definition, dso_local is set for ELF -fno-pic/-fpie and COFF and Mach-O. Since default clang -cc1 for ELF is similar to -fpic ("PIC Level" is not set), this nuance causes unneeded binary format differences. To make emitted IR similar, ELF -cc1 -fpic will default to -fno-semantic-interposition, which sets dso_local for default visibility external linkage definitions. To make this flip smooth and enable future (dso_local as definition default), this patch replaces (function) `define ` with `define{{.*}} `, (variable/constant/alias) `= ` with `={{.*}} `, or inserts appropriate `{{.*}} `.
2020-12-31 16:27:11 +08:00
// CHECK-DAG: define{{.*}} void @_Z16templated_kernelIiEvT_(
[CUDA] Make sure device-side __global__ functions are always visible. __global__ functions are a special case in CUDA. Even when the symbol would normally not be externally visible according to C++ rules, they still must be visible in CUDA GPU object so host-side stub can launch them. Differential Revision: http://reviews.llvm.org/D19748 llvm-svn: 268299
2016-05-03 04:30:03 +08:00
namespace {
__global__ void anonymous_ns_kernel() {}
[test] Add {{.*}} to make ELF tests immune to dso_local/dso_preemptable/(none) differences For a default visibility external linkage definition, dso_local is set for ELF -fno-pic/-fpie and COFF and Mach-O. Since default clang -cc1 for ELF is similar to -fpic ("PIC Level" is not set), this nuance causes unneeded binary format differences. To make emitted IR similar, ELF -cc1 -fpic will default to -fno-semantic-interposition, which sets dso_local for default visibility external linkage definitions. To make this flip smooth and enable future (dso_local as definition default), this patch replaces (function) `define ` with `define{{.*}} `, (variable/constant/alias) `= ` with `={{.*}} `, or inserts appropriate `{{.*}} `.
2020-12-31 16:27:11 +08:00
// CHECK-DAG: define{{.*}} void @_ZN12_GLOBAL__N_119anonymous_ns_kernelEv(
[CUDA] Make sure device-side __global__ functions are always visible. __global__ functions are a special case in CUDA. Even when the symbol would normally not be externally visible according to C++ rules, they still must be visible in CUDA GPU object so host-side stub can launch them. Differential Revision: http://reviews.llvm.org/D19748 llvm-svn: 268299
2016-05-03 04:30:03 +08:00
}
void host_function() {
templated_kernel<<<0, 0>>>(0);
anonymous_ns_kernel<<<0,0>>>();
}
[CUDA] Add implicit __attribute__((used)) to all __global__ functions. This makes sure that we emit kernels that were instantiated from the host code and which would never be explicitly referenced by anything else on device side. Differential Revision: http://reviews.llvm.org/D11666 llvm-svn: 248293
2015-09-23 01:22:51 +08:00
IR: Make metadata typeless in assembly, clang side Match LLVM changes from r224257. llvm-svn: 224259
2014-12-16 03:10:08 +08:00
// CHECK: !{{[0-9]+}} = !{void ()* @global_function, !"kernel", i32 1}
[CUDA] Add implicit __attribute__((used)) to all __global__ functions. This makes sure that we emit kernels that were instantiated from the host code and which would never be explicitly referenced by anything else on device side. Differential Revision: http://reviews.llvm.org/D11666 llvm-svn: 248293
2015-09-23 01:22:51 +08:00
// CHECK: !{{[0-9]+}} = !{void (i32)* @_Z16templated_kernelIiEvT_, !"kernel", i32 1}