65 lines · plain
1// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -ast-print -x hip -verify=NS,all %s2// RUN: %clang_cc1 -triple nvptx-nvidia-cuda -fcuda-is-device -ast-print -x hip -verify=NS,all %s3// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx1250 -fcuda-is-device -ast-print -x hip -verify=amd,common,all %s | FileCheck -check-prefixes=CHECK %s4// RUN: %clang_cc1 -triple nvptx-nvidia-cuda -target-cpu sm_90 -fcuda-is-device -ast-print -x hip -verify=cuda,common,all %s | FileCheck -check-prefixes=CHECK %s5// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -aux-triple amdgcn-amd-amdhsa -ast-print -x hip -verify=amd,common,all %s | FileCheck -check-prefixes=CHECK %s6 7#include "Inputs/cuda.h"8 9const int constint = 4;10 11// CHECK: __attribute__((global)) __attribute__((cluster_dims(2, 2, 2))) void test_literal_3d()12__global__ void __cluster_dims__(2, 2, 2) test_literal_3d() {} //NS-error {{'cluster_dims' is not supported for this GPU architecture}}13 14// CHECK: __attribute__((global)) __attribute__((cluster_dims(2, 2))) void test_literal_2d()15__global__ void __cluster_dims__(2, 2) test_literal_2d() {} //NS-error {{'cluster_dims' is not supported for this GPU architecture}}16 17// CHECK: __attribute__((global)) __attribute__((cluster_dims(4))) void test_literal_1d()18__global__ void __cluster_dims__(4) test_literal_1d() {} //NS-error {{'cluster_dims' is not supported for this GPU architecture}}19 20// CHECK: __attribute__((global)) __attribute__((cluster_dims(constint, constint / 4, 1))) void test_constant()21__global__ void __cluster_dims__(constint, constint / 4, 1) test_constant() {} //NS-error {{'cluster_dims' is not supported for this GPU architecture}}22 23// CHECK: template <int x, int y, int z> void test_template() __attribute__((cluster_dims(x, y, z)))24template <int x, int y, int z> void test_template(void) __cluster_dims__(x, y, z){} //NS-error {{'cluster_dims' is not supported for this GPU architecture}}25 26// CHECK: template <int x, int y, int z> void test_template_expr() __attribute__((cluster_dims(x + constint, y, z)))27template <int x, int y, int z> void test_template_expr(void) __cluster_dims__(x + constint, y, z) {} //NS-error {{'cluster_dims' is not supported for this GPU architecture}}28 29//NS-error@+1 {{'cluster_dims' is not supported for this GPU architecture}}30__global__ void __cluster_dims__(32, 2, 4) test_too_large_dim_0() {} // common-error {{integer constant expression evaluates to value 32 that cannot be represented in a 4-bit unsigned integer type}}31 32// cuda-error@+2 {{cluster does not support more than 8 thread blocks; 64 provided}}33// amd-error@+1 {{cluster does not support more than 16 thread blocks; 64 provided}}34__global__ void __cluster_dims__(4, 4, 4) test_too_large_dim_1() {} // NS-error {{'cluster_dims' is not supported for this GPU architecture}}35 36// cuda-error@+3 {{cluster does not support more than 8 thread blocks; 64 provided}}37// amd-error@+2 {{cluster does not support more than 16 thread blocks; 64 provided}}38template<unsigned a, unsigned b, unsigned c>39__global__ void __cluster_dims__(a, b, c) test_too_large_dim_template() {} // NS-error {{'cluster_dims' is not supported for this GPU architecture}}40template __global__ void test_too_large_dim_template<4, 4, 4>(); // common-note {{in instantiation of function template specialization 'test_too_large_dim_template<4U, 4U, 4U>' requested here}}41 42int none_const_int = 4;43 44//NS-error@+1 {{'cluster_dims' is not supported for this GPU architecture}}45__global__ void __cluster_dims__(none_const_int, 2, 4) test_non_constant_0() {} // common-error {{'cluster_dims' attribute requires parameter 0 to be an integer constant}}46 47//NS-error@+1 {{'cluster_dims' is not supported for this GPU architecture}}48__global__ void __cluster_dims__(8, none_const_int / 2, 4) test_non_constant_1() {} // common-error {{'cluster_dims' attribute requires parameter 1 to be an integer constant}}49 50//NS-error@+1 {{'cluster_dims' is not supported for this GPU architecture}}51__global__ void __cluster_dims__(8, 2, none_const_int / 4) test_non_constant_2() {} // common-error {{'cluster_dims' attribute requires parameter 2 to be an integer constant}}52 53//NS-error@+1 {{'no_cluster' is not supported for this GPU architecture}}54__global__ void __no_cluster__ test_no_cluster() {}55 56//NS-error@+2 {{'no_cluster' is not supported for this GPU architecture}}57//NS-error@+1 {{'cluster_dims' is not supported for this GPU architecture}}58__global__ void __no_cluster__ __cluster_dims__(2,2,2) test_have_both() {} // common-error {{'cluster_dims' and 'no_cluster' attributes are not compatible}} common-note {{conflicting attribute is here}}59 60template <int... args>61__cluster_dims__(args) void test_template_variadic_args(void) {} // all-error {{expression contains unexpanded parameter pack 'args'}}62 63template <int... args>64__cluster_dims__(1, args) void test_template_variadic_args_2(void) {} // all-error {{expression contains unexpanded parameter pack 'args'}}65