llvm-project/clang/test/SemaCUDA/spirv-attrs-diag.cu
Alexander Shaposhnikov 297f0b3f4c
[CudaSPIRV] Allow using integral non-type template parameters as attribute args (#131546)
Allow using integral non-type template parameters as attribute arguments
of
reqd_work_group_size and work_group_size_hint.

Test plan:
ninja check-all
2025-03-19 10:11:18 -07:00

59 lines
2.6 KiB
Plaintext

// RUN: %clang_cc1 -triple spirv64 -aux-triple x86_64-unknown-linux-gnu \
// RUN: -fcuda-is-device -verify -fsyntax-only %s
#define __global__ __attribute__((global))
__attribute__((reqd_work_group_size(0x100000000, 1, 1))) // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
__global__ void TestTooBigArg1(void);
__attribute__((work_group_size_hint(0x100000000, 1, 1))) // expected-error {{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}
__global__ void TestTooBigArg2(void);
template <int... Args>
__attribute__((reqd_work_group_size(Args))) // expected-error {{expression contains unexpanded parameter pack 'Args'}}
__global__ void TestTemplateVariadicArgs1(void) {}
template <int... Args>
__attribute__((work_group_size_hint(Args))) // expected-error {{expression contains unexpanded parameter pack 'Args'}}
__global__ void TestTemplateVariadicArgs2(void) {}
template <class a> // expected-note {{declared here}}
__attribute__((reqd_work_group_size(a, 1, 1))) // expected-error {{'a' does not refer to a value}}
__global__ void TestTemplateArgClass1(void) {}
template <class a> // expected-note {{declared here}}
__attribute__((work_group_size_hint(a, 1, 1))) // expected-error {{'a' does not refer to a value}}
__global__ void TestTemplateArgClass2(void) {}
constexpr int A = 512;
__attribute__((reqd_work_group_size(A, A, A)))
__global__ void TestConstIntArg1(void) {}
__attribute__((work_group_size_hint(A, A, A)))
__global__ void TestConstIntArg2(void) {}
int B = 512;
__attribute__((reqd_work_group_size(B, 1, 1))) // expected-error {{attribute requires parameter 0 to be an integer constant}}
__global__ void TestNonConstIntArg1(void) {}
__attribute__((work_group_size_hint(B, 1, 1))) // expected-error {{attribute requires parameter 0 to be an integer constant}}
__global__ void TestNonConstIntArg2(void) {}
constexpr int C = -512;
__attribute__((reqd_work_group_size(C, 1, 1))) // expected-error {{attribute requires a non-negative integral compile time constant expression}}
__global__ void TestNegativeConstIntArg1(void) {}
__attribute__((work_group_size_hint(C, 1, 1))) // expected-error {{attribute requires a non-negative integral compile time constant expression}}
__global__ void TestNegativeConstIntArg2(void) {}
__attribute__((reqd_work_group_size(A, 0, 1))) // expected-error {{attribute must be greater than 0}}
__global__ void TestZeroArg1(void) {}
__attribute__((work_group_size_hint(A, 0, 1))) // expected-error {{attribute must be greater than 0}}
__global__ void TestZeroArg2(void) {}