Artem Belevich | 7c3fdcc | 2024-11-05 10:48:54 -0800 | [diff] [blame] | 1 | // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --check-globals all --version 5 |
| 2 | // RUN: %clang_cc1 "-triple" "nvptx64-nvidia-cuda" -emit-llvm -fcuda-is-device -o - %s | FileCheck %s |
| 3 | |
| 4 | #include "Inputs/cuda.h" |
| 5 | |
| 6 | struct S {}; |
| 7 | |
| 8 | __global__ void kernel(__grid_constant__ const S gc_arg1, int arg2, __grid_constant__ const int gc_arg3) {} |
| 9 | |
| 10 | // dependent arguments get diagnosed after instantiation. |
| 11 | template <typename T> |
| 12 | __global__ void tkernel_const(__grid_constant__ const T arg) {} |
| 13 | |
| 14 | template <typename T> |
| 15 | __global__ void tkernel(int dummy, __grid_constant__ T arg) {} |
| 16 | |
| 17 | void foo() { |
| 18 | tkernel_const<const S><<<1,1>>>({}); |
| 19 | tkernel_const<S><<<1,1>>>({}); |
| 20 | tkernel<const S><<<1,1>>>(1, {}); |
| 21 | } |
| 22 | //. |
| 23 | //. |
Alex MacLean | 4583f6d | 2025-01-07 18:24:50 -0800 | [diff] [blame] | 24 | // CHECK: [[META0:![0-9]+]] = !{ptr @_Z6kernel1Sii, !"grid_constant", [[META1:![0-9]+]]} |
Artem Belevich | 7c3fdcc | 2024-11-05 10:48:54 -0800 | [diff] [blame] | 25 | // CHECK: [[META1]] = !{i32 1, i32 3} |
Alex MacLean | 4583f6d | 2025-01-07 18:24:50 -0800 | [diff] [blame] | 26 | // CHECK: [[META2:![0-9]+]] = !{ptr @_Z13tkernel_constIK1SEvT_, !"grid_constant", [[META3:![0-9]+]]} |
Artem Belevich | 7c3fdcc | 2024-11-05 10:48:54 -0800 | [diff] [blame] | 27 | // CHECK: [[META3]] = !{i32 1} |
Alex MacLean | 4583f6d | 2025-01-07 18:24:50 -0800 | [diff] [blame] | 28 | // CHECK: [[META4:![0-9]+]] = !{ptr @_Z13tkernel_constI1SEvT_, !"grid_constant", [[META3]]} |
| 29 | // CHECK: [[META5:![0-9]+]] = !{ptr @_Z7tkernelIK1SEviT_, !"grid_constant", [[META6:![0-9]+]]} |
Artem Belevich | 7c3fdcc | 2024-11-05 10:48:54 -0800 | [diff] [blame] | 30 | // CHECK: [[META6]] = !{i32 2} |
| 31 | //. |