329 lines · plain
1// RUN: %clang_cc1 -std=c++11 -fsyntax-only -verify %s2#include "Inputs/cuda.h"3 4 5__attribute__((amdgpu_flat_work_group_size(32, 64)))6__global__ void flat_work_group_size_32_64() {}7 8__attribute__((amdgpu_waves_per_eu(2)))9__global__ void waves_per_eu_2() {}10 11__attribute__((amdgpu_waves_per_eu(2, 4)))12__global__ void waves_per_eu_2_4() {}13 14__attribute__((amdgpu_num_sgpr(32)))15__global__ void num_sgpr_32() {}16 17__attribute__((amdgpu_num_vgpr(64)))18__global__ void num_vgpr_64() {}19 20 21__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2)))22__global__ void flat_work_group_size_32_64_waves_per_eu_2() {}23 24__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4)))25__global__ void flat_work_group_size_32_64_waves_per_eu_2_4() {}26 27__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_num_sgpr(32)))28__global__ void flat_work_group_size_32_64_num_sgpr_32() {}29 30__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_num_vgpr(64)))31__global__ void flat_work_group_size_32_64_num_vgpr_64() {}32 33__attribute__((amdgpu_waves_per_eu(2), amdgpu_num_sgpr(32)))34__global__ void waves_per_eu_2_num_sgpr_32() {}35 36__attribute__((amdgpu_waves_per_eu(2), amdgpu_num_vgpr(64)))37__global__ void waves_per_eu_2_num_vgpr_64() {}38 39__attribute__((amdgpu_waves_per_eu(2, 4), amdgpu_num_sgpr(32)))40__global__ void waves_per_eu_2_4_num_sgpr_32() {}41 42__attribute__((amdgpu_waves_per_eu(2, 4), amdgpu_num_vgpr(64)))43__global__ void waves_per_eu_2_4_num_vgpr_64() {}44 45__attribute__((amdgpu_num_sgpr(32), amdgpu_num_vgpr(64)))46__global__ void num_sgpr_32_num_vgpr_64() {}47 48__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2), amdgpu_num_sgpr(32)))49__global__ void flat_work_group_size_32_64_waves_per_eu_2_num_sgpr_32() {}50 51__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2), amdgpu_num_vgpr(64)))52__global__ void flat_work_group_size_32_64_waves_per_eu_2_num_vgpr_64() {}53 54__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4), amdgpu_num_sgpr(32)))55__global__ void flat_work_group_size_32_64_waves_per_eu_2_4_num_sgpr_32() {}56 57__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4), amdgpu_num_vgpr(64)))58__global__ void flat_work_group_size_32_64_waves_per_eu_2_4_num_vgpr_64() {}59 60__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2), amdgpu_num_sgpr(32), amdgpu_num_vgpr(64)))61__global__ void flat_work_group_size_32_64_waves_per_eu_2_num_sgpr_32_num_vgpr_64() {}62 63__attribute__((amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4), amdgpu_num_sgpr(32), amdgpu_num_vgpr(64)))64__global__ void flat_work_group_size_32_64_waves_per_eu_2_4_num_sgpr_32_num_vgpr_64() {}65 66__attribute__((amdgpu_max_num_work_groups(32, 1, 1)))67__global__ void max_num_work_groups_32_1_1() {}68 69__attribute__((amdgpu_max_num_work_groups(32, 1, 1), amdgpu_flat_work_group_size(32, 64)))70__global__ void max_num_work_groups_32_1_1_flat_work_group_size_32_64() {}71 72__attribute__((amdgpu_max_num_work_groups(32, 1, 1), amdgpu_flat_work_group_size(32, 64), amdgpu_waves_per_eu(2, 4), amdgpu_num_sgpr(32), amdgpu_num_vgpr(64)))73__global__ void max_num_work_groups_32_1_1_flat_work_group_size_32_64_waves_per_eu_2_4_num_sgpr_32_num_vgpr_64() {}74 75 76// expected-error@+2{{attribute 'reqd_work_group_size' can only be applied to an OpenCL kernel function}}77__attribute__((reqd_work_group_size(32, 64, 64)))78__global__ void reqd_work_group_size_32_64_64() {}79 80// expected-error@+2{{attribute 'work_group_size_hint' can only be applied to an OpenCL kernel function}}81__attribute__((work_group_size_hint(2, 2, 2)))82__global__ void work_group_size_hint_2_2_2() {}83 84// expected-error@+2{{attribute 'vec_type_hint' can only be applied to an OpenCL kernel function}}85__attribute__((vec_type_hint(int)))86__global__ void vec_type_hint_int() {}87 88// expected-error@+2{{attribute 'intel_reqd_sub_group_size' can only be applied to an OpenCL kernel function}}89__attribute__((intel_reqd_sub_group_size(64)))90__global__ void intel_reqd_sub_group_size_64() {}91 92// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be an integer constant}}93__attribute__((amdgpu_flat_work_group_size("32", 64)))94__global__ void non_int_min_flat_work_group_size_32_64() {}95// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 1 to be an integer constant}}96__attribute__((amdgpu_flat_work_group_size(32, "64")))97__global__ void non_int_max_flat_work_group_size_32_64() {}98 99int nc_min = 32, nc_max = 64;100// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be an integer constant}}101__attribute__((amdgpu_flat_work_group_size(nc_min, 64)))102__global__ void non_cint_min_flat_work_group_size_32_64() {}103// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 1 to be an integer constant}}104__attribute__((amdgpu_flat_work_group_size(32, nc_max)))105__global__ void non_cint_max_flat_work_group_size_32_64() {}106 107const int c_min = 16, c_max = 32;108__attribute__((amdgpu_flat_work_group_size(c_min * 2, 64)))109__global__ void cint_min_flat_work_group_size_32_64() {}110__attribute__((amdgpu_flat_work_group_size(32, c_max * 2)))111__global__ void cint_max_flat_work_group_size_32_64() {}112 113// expected-error@+3{{'T' does not refer to a value}}114// expected-note@+1{{declared here}}115template<typename T>116__attribute__((amdgpu_flat_work_group_size(T, 64)))117__global__ void template_class_min_flat_work_group_size_32_64() {}118// expected-error@+3{{'T' does not refer to a value}}119// expected-note@+1{{declared here}}120template<typename T>121__attribute__((amdgpu_flat_work_group_size(32, T)))122__global__ void template_class_max_flat_work_group_size_32_64() {}123 124template<unsigned a, unsigned b>125__attribute__((amdgpu_flat_work_group_size(a, b)))126__global__ void template_flat_work_group_size_32_64() {}127template __global__ void template_flat_work_group_size_32_64<32, 64>();128 129template<unsigned a, unsigned b, unsigned c>130__attribute__((amdgpu_flat_work_group_size(a + b, b + c)))131__global__ void template_complex_flat_work_group_size_32_64() {}132template __global__ void template_complex_flat_work_group_size_32_64<16, 16, 48>();133 134unsigned ipow2(unsigned n) { return n == 0 ? 1 : 2 * ipow2(n - 1); }135constexpr unsigned ce_ipow2(unsigned n) { return n == 0 ? 1 : 2 * ce_ipow2(n - 1); }136 137__attribute__((amdgpu_flat_work_group_size(ce_ipow2(5), ce_ipow2(6))))138__global__ void cexpr_flat_work_group_size_32_64() {}139// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 0 to be an integer constant}}140__attribute__((amdgpu_flat_work_group_size(ipow2(5), 64)))141__global__ void non_cexpr_min_flat_work_group_size_32_64() {}142// expected-error@+1{{'amdgpu_flat_work_group_size' attribute requires parameter 1 to be an integer constant}}143__attribute__((amdgpu_flat_work_group_size(32, ipow2(6))))144__global__ void non_cexpr_max_flat_work_group_size_32_64() {}145 146// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 0 to be an integer constant}}147__attribute__((amdgpu_waves_per_eu("2")))148__global__ void non_int_min_waves_per_eu_2() {}149// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 1 to be an integer constant}}150__attribute__((amdgpu_waves_per_eu(2, "4")))151__global__ void non_int_max_waves_per_eu_2_4() {}152 153// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 0 to be an integer constant}}154__attribute__((amdgpu_waves_per_eu(nc_min)))155__global__ void non_cint_min_waves_per_eu_2() {}156// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 1 to be an integer constant}}157__attribute__((amdgpu_waves_per_eu(2, nc_max)))158__global__ void non_cint_min_waves_per_eu_2_4() {}159 160__attribute__((amdgpu_waves_per_eu(c_min / 8)))161__global__ void cint_min_waves_per_eu_2() {}162__attribute__((amdgpu_waves_per_eu(c_min / 8, c_max / 8)))163__global__ void cint_min_waves_per_eu_2_4() {}164 165// expected-error@+3{{'T' does not refer to a value}}166// expected-note@+1{{declared here}}167template<typename T>168__attribute__((amdgpu_waves_per_eu(T)))169__global__ void cint_min_waves_per_eu_2() {}170// expected-error@+3{{'T' does not refer to a value}}171// expected-note@+1{{declared here}}172template<typename T>173__attribute__((amdgpu_waves_per_eu(2, T)))174__global__ void cint_min_waves_per_eu_2_4() {}175 176template<unsigned a>177__attribute__((amdgpu_waves_per_eu(a)))178__global__ void template_waves_per_eu_2() {}179template __global__ void template_waves_per_eu_2<2>();180 181template<unsigned a, unsigned b>182__attribute__((amdgpu_waves_per_eu(a, b)))183__global__ void template_waves_per_eu_2_4() {}184template __global__ void template_waves_per_eu_2_4<2, 4>();185 186template<unsigned a, unsigned b, unsigned c>187__attribute__((amdgpu_waves_per_eu(a + b, c - b)))188__global__ void template_complex_waves_per_eu_2_4() {}189template __global__ void template_complex_waves_per_eu_2_4<1, 1, 5>();190 191// expected-error@+2{{expression contains unexpanded parameter pack 'Args'}}192template<unsigned... Args>193__attribute__((amdgpu_waves_per_eu(Args)))194__global__ void template_waves_per_eu_2() {}195template __global__ void template_waves_per_eu_2<2, 4>();196 197__attribute__((amdgpu_waves_per_eu(ce_ipow2(1))))198__global__ void cexpr_waves_per_eu_2() {}199__attribute__((amdgpu_waves_per_eu(ce_ipow2(1), ce_ipow2(2))))200__global__ void cexpr_waves_per_eu_2_4() {}201// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 0 to be an integer constant}}202__attribute__((amdgpu_waves_per_eu(ipow2(1))))203__global__ void non_cexpr_waves_per_eu_2() {}204// expected-error@+1{{'amdgpu_waves_per_eu' attribute requires parameter 1 to be an integer constant}}205__attribute__((amdgpu_waves_per_eu(2, ipow2(2))))206__global__ void non_cexpr_waves_per_eu_2_4() {}207 208__attribute__((amdgpu_max_num_work_groups(32)))209__global__ void max_num_work_groups_32() {}210 211__attribute__((amdgpu_max_num_work_groups(32, 1)))212__global__ void max_num_work_groups_32_1() {}213 214// expected-error@+1{{'amdgpu_max_num_work_groups' attribute takes no more than 3 arguments}}215__attribute__((amdgpu_max_num_work_groups(32, 1, 1, 1)))216__global__ void max_num_work_groups_32_1_1_1() {}217 218// expected-error@+1{{'amdgpu_max_num_work_groups' attribute takes at least 1 argument}}219__attribute__((amdgpu_max_num_work_groups()))220__global__ void max_num_work_groups_no_arg() {}221 222// expected-error@+1{{expected expression}}223__attribute__((amdgpu_max_num_work_groups(,1,1)))224__global__ void max_num_work_groups_empty_1_1() {}225 226// expected-error@+1{{expected expression}}227__attribute__((amdgpu_max_num_work_groups(32,,1)))228__global__ void max_num_work_groups_32_empty_1() {}229 230// expected-error@+1{{'amdgpu_max_num_work_groups' attribute requires parameter 0 to be an integer constant}}231__attribute__((amdgpu_max_num_work_groups(ipow2(5), 1, 1)))232__global__ void max_num_work_groups_32_1_1_non_int_arg0() {}233 234// expected-error@+1{{'amdgpu_max_num_work_groups' attribute requires parameter 1 to be an integer constant}}235__attribute__((amdgpu_max_num_work_groups(32, "1", 1)))236__global__ void max_num_work_groups_32_1_1_non_int_arg1() {}237 238// expected-error@+1{{'amdgpu_max_num_work_groups' attribute requires a non-negative integral compile time constant expression}}239__attribute__((amdgpu_max_num_work_groups(-32, 1, 1)))240__global__ void max_num_work_groups_32_1_1_neg_int_arg0() {}241 242// expected-error@+1{{'amdgpu_max_num_work_groups' attribute requires a non-negative integral compile time constant expression}}243__attribute__((amdgpu_max_num_work_groups(32, -1, 1)))244__global__ void max_num_work_groups_32_1_1_neg_int_arg1() {}245 246// expected-error@+1{{'amdgpu_max_num_work_groups' attribute requires a non-negative integral compile time constant expression}}247__attribute__((amdgpu_max_num_work_groups(32, 1, -1)))248__global__ void max_num_work_groups_32_1_1_neg_int_arg2() {}249 250// expected-error@+1{{'amdgpu_max_num_work_groups' attribute must be greater than 0}}251__attribute__((amdgpu_max_num_work_groups(0, 1, 1)))252__global__ void max_num_work_groups_0_1_1() {}253 254// expected-error@+1{{'amdgpu_max_num_work_groups' attribute must be greater than 0}}255__attribute__((amdgpu_max_num_work_groups(32, 0, 1)))256__global__ void max_num_work_groups_32_0_1() {}257 258// expected-error@+1{{'amdgpu_max_num_work_groups' attribute must be greater than 0}}259__attribute__((amdgpu_max_num_work_groups(32, 1, 0)))260__global__ void max_num_work_groups_32_1_0() {}261 262__attribute__((amdgpu_max_num_work_groups(4294967295)))263__global__ void max_num_work_groups_max_unsigned_int() {}264 265// expected-error@+1{{integer constant expression evaluates to value 4294967296 that cannot be represented in a 32-bit unsigned integer type}}266__attribute__((amdgpu_max_num_work_groups(4294967296)))267__global__ void max_num_work_groups_max_unsigned_int_plus1() {}268 269// expected-error@+1{{integer constant expression evaluates to value 10000000000 that cannot be represented in a 32-bit unsigned integer type}}270__attribute__((amdgpu_max_num_work_groups(10000000000)))271__global__ void max_num_work_groups_too_large() {}272 273int num_wg_x = 32;274int num_wg_y = 1;275int num_wg_z = 1;276// expected-error@+1{{'amdgpu_max_num_work_groups' attribute requires parameter 0 to be an integer constant}}277__attribute__((amdgpu_max_num_work_groups(num_wg_x, 1, 1)))278__global__ void max_num_work_groups_32_1_1_non_const_arg0() {}279 280// expected-error@+1{{'amdgpu_max_num_work_groups' attribute requires parameter 1 to be an integer constant}}281__attribute__((amdgpu_max_num_work_groups(32, num_wg_y, 1)))282__global__ void max_num_work_groups_32_1_1_non_const_arg1() {}283 284// expected-error@+1{{'amdgpu_max_num_work_groups' attribute requires parameter 2 to be an integer constant}}285__attribute__((amdgpu_max_num_work_groups(32, 1, num_wg_z)))286__global__ void max_num_work_groups_32_1_1_non_const_arg2() {}287 288const int c_num_wg_x = 32;289__attribute__((amdgpu_max_num_work_groups(c_num_wg_x, 1, 1)))290__global__ void max_num_work_groups_32_1_1_const_arg0() {}291 292template<unsigned a>293__attribute__((amdgpu_max_num_work_groups(a, 1, 1)))294__global__ void template_a_1_1_max_num_work_groups() {}295template __global__ void template_a_1_1_max_num_work_groups<32>();296 297template<unsigned a>298__attribute__((amdgpu_max_num_work_groups(32, a, 1)))299__global__ void template_32_a_1_max_num_work_groups() {}300template __global__ void template_32_a_1_max_num_work_groups<1>();301 302template<unsigned a>303__attribute__((amdgpu_max_num_work_groups(32, 1, a)))304__global__ void template_32_1_a_max_num_work_groups() {}305template __global__ void template_32_1_a_max_num_work_groups<1>();306 307// expected-error@+3{{'amdgpu_max_num_work_groups' attribute must be greater than 0}}308// expected-note@+4{{in instantiation of}}309template<unsigned b>310__attribute__((amdgpu_max_num_work_groups(b, 1, 1)))311__global__ void template_b_1_1_max_num_work_groups() {}312template __global__ void template_b_1_1_max_num_work_groups<0>();313 314// expected-error@+3{{'amdgpu_max_num_work_groups' attribute must be greater than 0}}315// expected-note@+4{{in instantiation of}}316template<unsigned b>317__attribute__((amdgpu_max_num_work_groups(32, b, 1)))318__global__ void template_32_b_1_max_num_work_groups() {}319template __global__ void template_32_b_1_max_num_work_groups<0>();320 321// expected-error@+3{{'amdgpu_max_num_work_groups' attribute must be greater than 0}}322// expected-note@+4{{in instantiation of}}323template<unsigned b>324__attribute__((amdgpu_max_num_work_groups(32, 1, b)))325__global__ void template_32_1_b_max_num_work_groups() {}326template __global__ void template_32_1_b_max_num_work_groups<0>();327 328 329