diff --git a/clang/test/CodeGenSYCL/integration_header_ppmacros.cpp b/clang/test/CodeGenSYCL/integration_header_ppmacros.cpp index 752189ca5384..29b78a26b5a6 100644 --- a/clang/test/CodeGenSYCL/integration_header_ppmacros.cpp +++ b/clang/test/CodeGenSYCL/integration_header_ppmacros.cpp @@ -1,7 +1,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -sycl-std=2020 -fsycl-int-header=%t.h %s -// RUN: FileCheck -input-file=%t.h %s --check-prefix=CHECK-SYCL2020 -// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -sycl-std=2017 -fsycl-int-header=%t.h %s -// RUN: FileCheck -input-file=%t.h %s --check-prefix=CHECK-SYCL2017 +// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -fsycl-int-header=%t.h %s +// RUN: FileCheck -input-file=%t.h %s // RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -fsycl-range-rounding=disable -fsycl-int-header=%t.h %s // RUN: FileCheck -input-file=%t.h %s --check-prefix=CHECK-RANGE // RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -fsycl-range-rounding=force -fsycl-int-header=%t.h %s @@ -18,19 +16,9 @@ int main() { sycl::kernel_single_task([]() {}); } -// CHECK-SYCL2020: #ifndef SYCL_LANGUAGE_VERSION -// CHECK-SYCL2020-NEXT: #define SYCL_LANGUAGE_VERSION 202001 -// CHECK-SYCL2020-NEXT: #endif //SYCL_LANGUAGE_VERSION -// CHECK-SYCL2020-NOT: #define CL_SYCL_LANGUAGE_VERSION 121 -// CHECK-SYCL2020-NOT: #define SYCL_LANGUAGE_VERSION 201707 - -// CHECK-SYCL2017: #ifndef CL_SYCL_LANGUAGE_VERSION -// CHECK-SYCL2017-NEXT: #define CL_SYCL_LANGUAGE_VERSION 121 -// CHECK-SYCL2017-NEXT: #endif //CL_SYCL_LANGUAGE_VERSION -// CHECK-SYCL2017: #ifndef SYCL_LANGUAGE_VERSION -// CHECK-SYCL2017-NEXT: #define SYCL_LANGUAGE_VERSION 201707 -// CHECK-SYCL2017-NEXT: #endif //SYCL_LANGUAGE_VERSION -// CHECK-SYCL2017-NOT: #define SYCL_LANGUAGE_VERSION 202001 +// CHECK: #ifndef SYCL_LANGUAGE_VERSION +// CHECK-NEXT: #define SYCL_LANGUAGE_VERSION 202001 +// CHECK-NEXT: #endif //SYCL_LANGUAGE_VERSION // CHECK-RANGE: #ifndef __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ // CHECK-RANGE-NEXT: #define __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ 1 diff --git a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp index 17a9d764b08d..9a391a6cbba9 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" @@ -16,9 +16,6 @@ class Functor { [[intel::no_global_work_offset(SIZE)]] void operator()() const {} }; -template -[[intel::no_global_work_offset(N)]] void func() {} - int main() { q.submit([&](handler &h) { Foo boo; @@ -32,10 +29,6 @@ int main() { Functor<1> f; h.single_task(f); - - h.single_task([]() { - func<1>(); - }); }); return 0; } @@ -44,6 +37,5 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !no_global_work_offset ![[NUM5]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} ![[NUM4:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !no_global_work_offset ![[NUM5]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !no_global_work_offset ![[NUM5]] // CHECK-NOT: ![[NUM4]] = !{i32 0} // CHECK: ![[NUM5]] = !{} diff --git a/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp b/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp index c6f7baf27054..392d5fe219c0 100644 --- a/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp +++ b/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" @@ -16,9 +16,6 @@ class Functor { [[intel::max_global_work_dim(SIZE)]] void operator()() const {} }; -template -[[intel::max_global_work_dim(N)]] void func() {} - int main() { q.submit([&](handler &h) { Foo boo; @@ -29,10 +26,6 @@ int main() { Functor<2> f; h.single_task(f); - - h.single_task([]() { - func<2>(); - }); }); return 0; } @@ -40,6 +33,5 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !max_global_work_dim ![[NUM1:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !max_global_work_dim ![[NUM2:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !max_global_work_dim ![[NUM2]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !max_global_work_dim ![[NUM2]] // CHECK: ![[NUM1]] = !{i32 1} // CHECK: ![[NUM2]] = !{i32 2} diff --git a/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp b/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp index 97ab18797465..76b47b3ffc8f 100644 --- a/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp +++ b/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" @@ -21,8 +21,6 @@ class Functor { [[intel::max_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() const {} }; -template -[[intel::max_work_group_size(N, N1, N2)]] void func() {} int main() { q.submit([&](handler &h) { @@ -38,9 +36,6 @@ int main() { Functor<2, 2, 2> f; h.single_task(f); - h.single_task([]() { - func<4, 4, 4>(); - }); }); return 0; } @@ -49,9 +44,7 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !max_work_group_size ![[NUM8:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !max_work_group_size ![[NUM6:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !max_work_group_size ![[NUM2:[0-9]+]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !max_work_group_size ![[NUM4:[0-9]+]] // CHECK: ![[NUM1]] = !{i32 1, i32 1, i32 1} // CHECK: ![[NUM8]] = !{i32 8, i32 8, i32 8} // CHECK: ![[NUM6]] = !{i32 6, i32 3, i32 1} // CHECK: ![[NUM2]] = !{i32 2, i32 2, i32 2} -// CHECK: ![[NUM4]] = !{i32 4, i32 4, i32 4} diff --git a/clang/test/CodeGenSYCL/intel-restrict.cpp b/clang/test/CodeGenSYCL/intel-restrict.cpp index 9845da905e0e..e25abdcf1447 100644 --- a/clang/test/CodeGenSYCL/intel-restrict.cpp +++ b/clang/test/CodeGenSYCL/intel-restrict.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device %s -emit-llvm -sycl-std=2017 -triple spir64-unknown-unknown -o - | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device %s -emit-llvm -triple spir64-unknown-unknown -o - | FileCheck %s template __attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { diff --git a/clang/test/CodeGenSYCL/kernel-by-reference.cpp b/clang/test/CodeGenSYCL/kernel-by-reference.cpp index cfc2cd61c490..d5506c6a629b 100644 --- a/clang/test/CodeGenSYCL/kernel-by-reference.cpp +++ b/clang/test/CodeGenSYCL/kernel-by-reference.cpp @@ -1,12 +1,12 @@ -// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -sycl-std=2017 -DSYCL2017 %s -// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -sycl-std=2020 -DSYCL2020 %s -// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -Wno-sycl-strict -DNODIAG %s -// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -sycl-std=2020 -Wno-sycl-strict -DNODIAG %s +// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only %s +// RUN: %clang_cc1 -triple spir64 -fsycl-is-device -internal-isystem %S/Inputs -verify -fsyntax-only -Wno-sycl-strict %s #include "sycl.hpp" using namespace sycl; +// expected-no-diagnostics + int simple_add(int i) { return i + 1; } @@ -14,24 +14,10 @@ int simple_add(int i) { // ensure both compile. int main() { queue q; -#if defined(SYCL2020) - // expected-warning@#KernelSingleTask2017 {{passing kernel functions by value is deprecated in SYCL 2020}} - // expected-note@+3 {{in instantiation of function template specialization}} -#endif - q.submit([&](handler &h) { - h.single_task_2017([]() { simple_add(10); }); - }); -#if defined(SYCL2017) - // expected-warning@#KernelSingleTask {{passing of kernel functions by reference is a SYCL 2020 extension}} - // expected-note@+3 {{in instantiation of function template specialization}} -#endif q.submit([&](handler &h) { h.single_task([]() { simple_add(11); }); }); return 0; } -#if defined(NODIAG) -// expected-no-diagnostics -#endif diff --git a/clang/test/CodeGenSYCL/launch_bounds_nvptx.cpp b/clang/test/CodeGenSYCL/launch_bounds_nvptx.cpp index 879fd79c1fe5..f85f2908d468 100644 --- a/clang/test/CodeGenSYCL/launch_bounds_nvptx.cpp +++ b/clang/test/CodeGenSYCL/launch_bounds_nvptx.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple nvptx-unknown-unknown -target-cpu sm_90 -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple nvptx-unknown-unknown -target-cpu sm_90 -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s // Test correct handling of maximum work group size, minimum work groups per // compute unit and maximum work groups per multi-processor attributes, that @@ -25,15 +25,6 @@ template class Functor { operator()() const {} }; -template -[[intel::max_work_group_size(N, 4, 8), intel::min_work_groups_per_cu(N), - intel::max_work_groups_per_mp(N)]] void -zoo() {} - -[[intel::max_work_group_size(2, 4, 8), intel::min_work_groups_per_cu(2), - intel::max_work_groups_per_mp(4)]] void -bar() {} - int main() { q.submit([&](handler &h) { // Test attribute argument size. @@ -49,12 +40,6 @@ int main() { // Test class template argument. Functor<6> f; h.single_task(f); - - // Test attribute is propagated. - h.single_task([]() { bar(); }); - - // Test function template argument. - h.single_task([]() { zoo<16>(); }); }); return 0; } @@ -62,8 +47,6 @@ int main() { // CHECK: define dso_local void @{{.*}}kernel_name1() #0 {{.*}} !min_work_groups_per_cu ![[MWGPC:[0-9]+]] !max_work_groups_per_mp ![[MWGPM:[0-9]+]] !max_work_group_size ![[MWGS:[0-9]+]] // CHECK: define dso_local void @{{.*}}kernel_name2() #0 {{.*}} !min_work_groups_per_cu ![[MWGPC:[0-9]+]] !max_work_groups_per_mp ![[MWGPM:[0-9]+]] !max_work_group_size ![[MWGS:[0-9]+]] // CHECK: define dso_local void @{{.*}}kernel_name3() #0 {{.*}} !min_work_groups_per_cu ![[MWGPC_MWGPM:[0-9]+]] !max_work_groups_per_mp ![[MWGPC_MWGPM]] !max_work_group_size ![[MWGS_2:[0-9]+]] -// CHECK: define dso_local void @{{.*}}kernel_name4() #0 {{.*}} !min_work_groups_per_cu ![[MWGPC:[0-9]+]] !max_work_groups_per_mp ![[MWGPM:[0-9]+]] !max_work_group_size ![[MWGS:[0-9]+]] -// CHECK: define dso_local void @{{.*}}kernel_name5() #0 {{.*}} !min_work_groups_per_cu ![[MWGPC_MWGPM_2:[0-9]+]] !max_work_groups_per_mp ![[MWGPC_MWGPM_2]] !max_work_group_size ![[MWGS_3:[0-9]+]] // CHECK: {{.*}}@{{.*}}kernel_name1, !"maxntidx", i32 8} // CHECK: {{.*}}@{{.*}}kernel_name1, !"maxntidy", i32 4} @@ -95,31 +78,9 @@ int main() { // CHECK: {{.*}}@{{.*}}Functor{{.*}}, !"maxntidz", i32 6} // CHECK: {{.*}}@{{.*}}Functor{{.*}}, !"minctasm", i32 6} // CHECK: {{.*}}@{{.*}}Functor{{.*}}, !"maxclusterrank", i32 6} -// CHECK: {{.*}}@{{.*}}kernel_name4, !"maxntidx", i32 8} -// CHECK: {{.*}}@{{.*}}kernel_name4, !"maxntidy", i32 4} -// CHECK: {{.*}}@{{.*}}kernel_name4, !"maxntidz", i32 2} -// CHECK: {{.*}}@{{.*}}kernel_name4, !"minctasm", i32 2} -// CHECK: {{.*}}@{{.*}}kernel_name4, !"maxclusterrank", i32 4} -// CHECK: {{.*}}@{{.*}}bar{{.*}}, !"maxntidx", i32 8} -// CHECK: {{.*}}@{{.*}}bar{{.*}}, !"maxntidy", i32 4} -// CHECK: {{.*}}@{{.*}}bar{{.*}}, !"maxntidz", i32 2} -// CHECK: {{.*}}@{{.*}}bar{{.*}}, !"minctasm", i32 2} -// CHECK: {{.*}}@{{.*}}bar{{.*}}, !"maxclusterrank", i32 4} -// CHECK: {{.*}}@{{.*}}kernel_name5, !"maxntidx", i32 8} -// CHECK: {{.*}}@{{.*}}kernel_name5, !"maxntidy", i32 4} -// CHECK: {{.*}}@{{.*}}kernel_name5, !"maxntidz", i32 16} -// CHECK: {{.*}}@{{.*}}kernel_name5, !"minctasm", i32 16} -// CHECK: {{.*}}@{{.*}}kernel_name5, !"maxclusterrank", i32 16} -// CHECK: {{.*}}@{{.*}}zoo{{.*}}, !"maxntidx", i32 8} -// CHECK: {{.*}}@{{.*}}zoo{{.*}}, !"maxntidy", i32 4} -// CHECK: {{.*}}@{{.*}}zoo{{.*}}, !"maxntidz", i32 16} -// CHECK: {{.*}}@{{.*}}zoo{{.*}}, !"minctasm", i32 16} -// CHECK: {{.*}}@{{.*}}zoo{{.*}}, !"maxclusterrank", i32 16} // CHECK: ![[MWGPC]] = !{i32 2} // CHECK: ![[MWGPM]] = !{i32 4} // CHECK: ![[MWGS]] = !{i32 8, i32 4, i32 2} // CHECK: ![[MWGPC_MWGPM]] = !{i32 6} // CHECK: ![[MWGS_2]] = !{i32 8, i32 4, i32 6} -// CHECK: ![[MWGPC_MWGPM_2]] = !{i32 16} -// CHECK: ![[MWGS_3]] = !{i32 8, i32 4, i32 16} diff --git a/clang/test/CodeGenSYCL/num-simd-work-items.cpp b/clang/test/CodeGenSYCL/num-simd-work-items.cpp index aeba63e0ffee..e613510f11b4 100644 --- a/clang/test/CodeGenSYCL/num-simd-work-items.cpp +++ b/clang/test/CodeGenSYCL/num-simd-work-items.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" @@ -16,8 +16,6 @@ class Functor { [[intel::num_simd_work_items(SIZE)]] void operator()() const {} }; -template -[[intel::num_simd_work_items(N)]] void func() {} int main() { q.submit([&](handler &h) { @@ -30,9 +28,6 @@ int main() { Functor<2> f; h.single_task(f); - h.single_task([]() { - func<4>(); - }); }); return 0; } @@ -40,8 +35,6 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !num_simd_work_items ![[NUM1:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !num_simd_work_items ![[NUM42:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !num_simd_work_items ![[NUM2:[0-9]+]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !num_simd_work_items ![[NUM4:[0-9]+]] // CHECK: ![[NUM1]] = !{i32 1} // CHECK: ![[NUM42]] = !{i32 42} // CHECK: ![[NUM2]] = !{i32 2} -// CHECK: ![[NUM4]] = !{i32 4} diff --git a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp index 6df63cea4dc0..4396313eac48 100644 --- a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp +++ b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -sycl-std=2017 -triple spir64-unknown-unknown -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -triple spir64-unknown-unknown -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" @@ -10,15 +10,6 @@ class Functor16 { [[intel::reqd_sub_group_size(16)]] void operator()() const {} }; -[[intel::reqd_sub_group_size(8)]] void foo() {} - -class Functor8 { -public: - void operator()() const { - foo(); - } -}; - template class Functor2 { public: @@ -33,9 +24,6 @@ int main() { Functor16 f16; h.single_task(f16); - Functor8 f8; - h.single_task(f8); - h.single_task( []() [[intel::reqd_sub_group_size(4)]]{}); @@ -50,11 +38,8 @@ int main() { } // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE16:[0-9]+]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE8:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE4:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE2:[0-9]+]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE2]] // CHECK: ![[SGSIZE16]] = !{i32 16} -// CHECK: ![[SGSIZE8]] = !{i32 8} // CHECK: ![[SGSIZE4]] = !{i32 4} // CHECK: ![[SGSIZE2]] = !{i32 2} diff --git a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp index 2eff74dc8776..542655a94ac3 100644 --- a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp +++ b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp @@ -1,7 +1,7 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple amdgcn-amd-amdhsa -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple nvptx-nvidia-cuda -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple nvptx64-nvidia-cuda -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple amdgcn-amd-amdhsa -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple nvptx-nvidia-cuda -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple nvptx64-nvidia-cuda -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" @@ -25,35 +25,6 @@ class CLFunctor32x16x16 { [[cl::reqd_work_group_size(32, 16, 16)]] void operator()() const {} }; -[[sycl::reqd_work_group_size(8)]] void f8() {} -[[sycl::reqd_work_group_size(8, 1)]] void f8x1() {} -[[sycl::reqd_work_group_size(8, 1, 1)]] void f8x1x1() {} -[[cl::reqd_work_group_size(8, 1, 1)]] void clf8x1x1() {} - -class Functor1D { -public: - void operator()() const { - f8(); - } -}; -class Functor2D { -public: - void operator()() const { - f8x1(); - } -}; -class Functor3D { -public: - void operator()() const { - f8x1x1(); - } -}; -class CLFunctor3D { -public: - void operator()() const { - clf8x1x1(); - } -}; template class FunctorTemp1D { @@ -76,41 +47,24 @@ class CLFunctorTemp3D { [[cl::reqd_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() const {} }; -template -[[sycl::reqd_work_group_size(N)]] void func1D() {} -template -[[sycl::reqd_work_group_size(N, N1)]] void func2D() {} -template -[[sycl::reqd_work_group_size(N, N1, N2)]] void func3D() {} -template -[[cl::reqd_work_group_size(N, N1, N2)]] void clfunc3D() {} - int main() { q.submit([&](handler &h) { CLFunctor32x16x16 clf32x16x16; h.single_task(clf32x16x16); - CLFunctor3D clf3d; - h.single_task(clf3d); - + h.single_task( []() [[cl::reqd_work_group_size(8, 8, 8)]]{}); CLFunctorTemp3D<2, 2, 2> clft3d; h.single_task(clft3d); - h.single_task([]() { - clfunc3D<8, 4, 4>(); - }); - h.single_task( []() [[cl::reqd_work_group_size(1, 8, 2)]]{}); Functor32x16x16 f32x16x16; h.single_task(f32x16x16); - Functor3D f3d; - h.single_task(f3d); h.single_task( []() [[sycl::reqd_work_group_size(8, 8, 8)]]{}); @@ -118,9 +72,6 @@ int main() { FunctorTemp3D<2, 2, 2> ft3d; h.single_task(ft3d); - h.single_task([]() { - func3D<8, 4, 4>(); - }); h.single_task( []() [[sycl::reqd_work_group_size(1, 8, 2)]]{}); @@ -128,8 +79,6 @@ int main() { Functor32x16 f32x16; h.single_task(f32x16); - Functor2D f2d; - h.single_task(f2d); h.single_task( []() [[sycl::reqd_work_group_size(8, 8)]]{}); @@ -137,9 +86,6 @@ int main() { FunctorTemp2D<2, 2> ft2d; h.single_task(ft2d); - h.single_task([]() { - func2D<8, 4>(); - }); h.single_task( []() [[sycl::reqd_work_group_size(1, 8)]]{}); @@ -147,8 +93,6 @@ int main() { Functor32 f32; h.single_task(f32); - Functor1D f1d; - h.single_task(f1d); h.single_task( []() [[sycl::reqd_work_group_size(8)]]{}); @@ -156,10 +100,6 @@ int main() { FunctorTemp1D<2> ft1d; h.single_task(ft1d); - h.single_task([]() { - func1D<8>(); - }); - h.single_task( []() [[sycl::reqd_work_group_size(1)]]{}); }); @@ -167,43 +107,31 @@ int main() { } // CHECK: define {{.*}} void @{{.*}}kernel_name1() #0 {{.*}} !work_group_num_dim ![[NDRWGS3D:[0-9]+]] !reqd_work_group_size ![[WGSIZE3D32:[0-9]+]] -// CHECK: define {{.*}} void @{{.*}}kernel_name2() #0 {{.*}} !work_group_num_dim ![[NDRWGS3D:[0-9]+]] !reqd_work_group_size ![[WGSIZE3D8:[0-9]+]] // CHECK: define {{.*}} void @{{.*}}kernel_name3() #0 {{.*}} !work_group_num_dim ![[NDRWGS3D:[0-9]+]] !reqd_work_group_size ![[WGSIZE3D88:[0-9]+]] // CHECK: define {{.*}} void @{{.*}}kernel_name4() #0 {{.*}} !work_group_num_dim ![[NDRWGS3D:[0-9]+]] !reqd_work_group_size ![[WGSIZE3D22:[0-9]+]] -// CHECK: define {{.*}} void @{{.*}}kernel_name5() #0 {{.*}} !work_group_num_dim ![[NDRWGS3D:[0-9]+]] !reqd_work_group_size ![[WGSIZE3D44:[0-9]+]] // CHECK: define {{.*}} void @{{.*}}kernel_name6() #0 {{.*}} !work_group_num_dim ![[NDRWGS3D:[0-9]+]] !reqd_work_group_size ![[WGSIZE3D2:[0-9]+]] // CHECK: define {{.*}} void @{{.*}}kernel_name7() #0 {{.*}} !work_group_num_dim ![[NDRWGS3D:[0-9]+]] !reqd_work_group_size ![[WGSIZE3D32]] -// CHECK: define {{.*}} void @{{.*}}kernel_name8() #0 {{.*}} !work_group_num_dim ![[NDRWGS3D:[0-9]+]] !reqd_work_group_size ![[WGSIZE3D8]] // CHECK: define {{.*}} void @{{.*}}kernel_name9() #0 {{.*}} !work_group_num_dim ![[NDRWGS3D:[0-9]+]] !reqd_work_group_size ![[WGSIZE3D88]] // CHECK: define {{.*}} void @{{.*}}kernel_name10() #0 {{.*}} !work_group_num_dim ![[NDRWGS3D:[0-9]+]] !reqd_work_group_size ![[WGSIZE3D22]] -// CHECK: define {{.*}} void @{{.*}}kernel_name11() #0 {{.*}} !work_group_num_dim ![[NDRWGS3D:[0-9]+]] !reqd_work_group_size ![[WGSIZE3D44]] // CHECK: define {{.*}} void @{{.*}}kernel_name12() #0 {{.*}} !work_group_num_dim ![[NDRWGS3D:[0-9]+]] !reqd_work_group_size ![[WGSIZE3D2]] // CHECK: define {{.*}} void @{{.*}}kernel_name13() #0 {{.*}} !work_group_num_dim ![[NDRWGS2D:[0-9]+]] !reqd_work_group_size ![[WGSIZE2D32:[0-9]+]] -// CHECK: define {{.*}} void @{{.*}}kernel_name14() #0 {{.*}} !work_group_num_dim ![[NDRWGS2D:[0-9]+]] !reqd_work_group_size ![[WGSIZE2D8:[0-9]+]] // CHECK: define {{.*}} void @{{.*}}kernel_name15() #0 {{.*}} !work_group_num_dim ![[NDRWGS2D:[0-9]+]] !reqd_work_group_size ![[WGSIZE2D88:[0-9]+]] // CHECK: define {{.*}} void @{{.*}}kernel_name16() #0 {{.*}} !work_group_num_dim ![[NDRWGS2D:[0-9]+]] !reqd_work_group_size ![[WGSIZE2D22:[0-9]+]] -// CHECK: define {{.*}} void @{{.*}}kernel_name17() #0 {{.*}} !work_group_num_dim ![[NDRWGS2D:[0-9]+]] !reqd_work_group_size ![[WGSIZE2D44:[0-9]+]] // CHECK: define {{.*}} void @{{.*}}kernel_name18() #0 {{.*}} !work_group_num_dim ![[NDRWGS2D:[0-9]+]] !reqd_work_group_size ![[WGSIZE2D2_or_WGSIZE1D8:[0-9]+]] // CHECK: define {{.*}} void @{{.*}}kernel_name19() #0 {{.*}} !work_group_num_dim ![[NDRWGS1D:[0-9]+]] !reqd_work_group_size ![[WGSIZE1D32:[0-9]+]] -// CHECK: define {{.*}} void @{{.*}}kernel_name20() #0 {{.*}} !work_group_num_dim ![[NDRWGS1D:[0-9]+]] !reqd_work_group_size ![[WGSIZE2D2_or_WGSIZE1D8]] // CHECK: define {{.*}} void @{{.*}}kernel_name21() #0 {{.*}} !work_group_num_dim ![[NDRWGS1D:[0-9]+]] !reqd_work_group_size ![[WGSIZE2D2_or_WGSIZE1D8]] // CHECK: define {{.*}} void @{{.*}}kernel_name22() #0 {{.*}} !work_group_num_dim ![[NDRWGS1D:[0-9]+]] !reqd_work_group_size ![[WGSIZE1D22:[0-9]+]] -// CHECK: define {{.*}} void @{{.*}}kernel_name23() #0 {{.*}} !work_group_num_dim ![[NDRWGS1D:[0-9]+]] !reqd_work_group_size ![[WGSIZE2D2_or_WGSIZE1D8]] // CHECK: define {{.*}} void @{{.*}}kernel_name24() #0 {{.*}} !work_group_num_dim ![[NDRWGS1D:[0-9]+]] !reqd_work_group_size ![[WGSIZE1D2:[0-9]+]] // CHECK: ![[NDRWGS3D]] = !{i32 3} // CHECK: ![[WGSIZE3D32]] = !{i32 16, i32 16, i32 32} -// CHECK: ![[WGSIZE3D8]] = !{i32 1, i32 1, i32 8} // CHECK: ![[WGSIZE3D88]] = !{i32 8, i32 8, i32 8} // CHECK: ![[WGSIZE3D22]] = !{i32 2, i32 2, i32 2} -// CHECK: ![[WGSIZE3D44]] = !{i32 4, i32 4, i32 8} // CHECK: ![[WGSIZE3D2]] = !{i32 2, i32 8, i32 1} // CHECK: ![[NDRWGS2D]] = !{i32 2} // CHECK: ![[WGSIZE2D32]] = !{i32 16, i32 32, i32 1} -// CHECK: ![[WGSIZE2D8]] = !{i32 1, i32 8, i32 1} // CHECK: ![[WGSIZE2D88]] = !{i32 8, i32 8, i32 1} // CHECK: ![[WGSIZE2D22]] = !{i32 2, i32 2, i32 1} -// CHECK: ![[WGSIZE2D44]] = !{i32 4, i32 8, i32 1} // CHECK: ![[WGSIZE2D2_or_WGSIZE1D8]] = !{i32 8, i32 1, i32 1} // CHECK: ![[NDRWGS1D]] = !{i32 1} // CHECK: ![[WGSIZE1D32]] = !{i32 32, i32 1, i32 1} diff --git a/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp b/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp index 5194b92f4458..4b3c2cbdea99 100644 --- a/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp +++ b/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" @@ -16,11 +16,6 @@ class Functor { [[intel::scheduler_target_fmax_mhz(N)]] void operator()() const {} }; -template -[[intel::scheduler_target_fmax_mhz(N)]] void zoo() {} - -[[intel::scheduler_target_fmax_mhz(2)]] void bar() {} - int main() { q.submit([&](handler &h) { // Test attribute argument size. @@ -35,13 +30,6 @@ int main() { Functor<7> f; h.single_task(f); - // Test attribute is propagated. - h.single_task( - []() { bar(); }); - - // Test function template argument. - h.single_task( - []() { zoo<75>(); }); }); return 0; } @@ -49,10 +37,6 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM5:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM42:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM7:[0-9]+]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM2:[0-9]+]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM75:[0-9]+]] // CHECK: ![[NUM5]] = !{i32 5} // CHECK: ![[NUM42]] = !{i32 42} // CHECK: ![[NUM7]] = !{i32 7} -// CHECK: ![[NUM2]] = !{i32 2} -// CHECK: ![[NUM75]] = !{i32 75} diff --git a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp index 0e727251a1ce..45f258a01b67 100644 --- a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp +++ b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" @@ -21,9 +21,6 @@ class Functor2 { [[sycl::reqd_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() const {} }; -template -[[sycl::reqd_work_group_size(N, N1, N2)]] void func() {} - int main() { q.submit([&](handler &h) { Functor foo; @@ -34,10 +31,6 @@ int main() { Functor2<2, 2, 2> foo2; h.single_task(foo2); - - h.single_task([]() { - func<8, 4, 4>(); - }); }); return 0; } @@ -45,10 +38,8 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !reqd_work_group_size ![[WGSIZE:[0-9]+]] !intel_reqd_sub_group_size ![[SGSIZE:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !reqd_work_group_size ![[WGSIZE1:[0-9]+]] !intel_reqd_sub_group_size ![[SGSIZE1:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !reqd_work_group_size ![[WGSIZE2:[0-9]+]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !reqd_work_group_size ![[WGSIZE3:[0-9]+]] // CHECK: ![[WGSIZE]] = !{i32 16, i32 16, i32 32} // CHECK: ![[SGSIZE]] = !{i32 4} // CHECK: ![[WGSIZE1]] = !{i32 32, i32 32, i32 64} // CHECK: ![[SGSIZE1]] = !{i32 2} // CHECK: ![[WGSIZE2]] = !{i32 2, i32 2, i32 2} -// CHECK: ![[WGSIZE3]] = !{i32 4, i32 4, i32 8} diff --git a/clang/test/Driver/sycl-device.cpp b/clang/test/Driver/sycl-device.cpp index 8d806ef5a7a7..5af8b5011161 100644 --- a/clang/test/Driver/sycl-device.cpp +++ b/clang/test/Driver/sycl-device.cpp @@ -18,7 +18,7 @@ // RUN: | FileCheck -check-prefix=CHECK-SYCL-NO_STRICT %s // CHECK-SYCL-NO_STRICT: clang{{.*}} "-Wno-sycl-strict" -/// Check that -sycl-std=2017 is set if no std version is provided by user +/// Check that -sycl-std=2020 is set if no std version is provided by user // RUN: %clang -### -fsycl %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHECK-SYCL-STD_VERSION %s // CHECK-SYCL-STD_VERSION: clang{{.*}} "-sycl-std=2020" diff --git a/clang/test/Driver/sycl.c b/clang/test/Driver/sycl.c index 05894c8811ae..a330247cad42 100644 --- a/clang/test/Driver/sycl.c +++ b/clang/test/Driver/sycl.c @@ -1,18 +1,12 @@ // RUN: %clang -### -fsycl -c %s 2>&1 | FileCheck %s --check-prefix=ENABLED // RUN: %clang -### -fsycl %s 2>&1 | FileCheck %s --check-prefix=ENABLED -// RUN: %clang -### -fsycl -sycl-std=1.2.1 %s 2>&1 | FileCheck %s --check-prefix=ENABLED -// RUN: %clang -### -fsycl -sycl-std=121 %s 2>&1 | FileCheck %s --check-prefix=ENABLED -// RUN: %clang -### -fsycl -sycl-std=2017 %s 2>&1 | FileCheck %s --check-prefix=ENABLED -// RUN: %clang -### -fsycl -sycl-std=2020 %s 2>&1 | FileCheck %s --check-prefix=ENABLED -// RUN: %clang -### -fsycl -sycl-std=sycl-1.2.1 %s 2>&1 | FileCheck %s --check-prefix=ENABLED // RUN: %clang -### -fno-sycl -fsycl %s 2>&1 | FileCheck %s --check-prefix=ENABLED -// RUN: %clang -### -sycl-std=2017 %s 2>&1 | FileCheck %s --check-prefix=NOT_ENABLED +// RUN: %clang -### %s 2>&1 | FileCheck %s --check-prefix=NOT_ENABLED // RUN: %clangxx -### -fsycl %s 2>&1 | FileCheck %s --check-prefix=ENABLED // RUN: %clangxx -### -fno-sycl %s 2>&1 | FileCheck %s --check-prefix=DISABLED // RUN: %clangxx -### -fsycl -fno-sycl %s 2>&1 | FileCheck %s --check-prefix=DISABLED // RUN: %clangxx -### %s 2>&1 | FileCheck %s --check-prefix=NOT_ENABLED // RUN: %clangxx -### -sycl-std=some-std %s 2>&1 | FileCheck %s --check-prefix=NOT_ENABLED -// RUN: %clang_cl -### -fsycl -sycl-std=2017 -- %s 2>&1 | FileCheck %s --check-prefix=ENABLED // RUN: %clang_cl -### -fsycl -- %s 2>&1 | FileCheck %s --check-prefix=ENABLED // RUN: %clang_cl -### -- %s 2>&1 | FileCheck %s --check-prefix=NOT_ENABLED // RUN: %clang_cl -### -sycl-std=some-std -- %s 2>&1 | FileCheck %s --check-prefix=NOT_ENABLED diff --git a/clang/test/Frontend/sycl-unnamed-lambdas-default.cpp b/clang/test/Frontend/sycl-unnamed-lambdas-default.cpp index 6faec626dadd..47d4c33dba73 100644 --- a/clang/test/Frontend/sycl-unnamed-lambdas-default.cpp +++ b/clang/test/Frontend/sycl-unnamed-lambdas-default.cpp @@ -1,18 +1,17 @@ // RUN: %clang_cc1 -fsycl-is-device %s -verify -DUNNAMED_LAMBDAS // RUN: %clang_cc1 -fsycl-is-host %s -verify -DUNNAMED_LAMBDAS // -// RUN: %clang_cc1 -fsycl-is-host -sycl-std=2017 %s -verify -// RUN: %clang_cc1 -fsycl-is-host -sycl-std=2020 %s -verify -DUNNAMED_LAMBDAS +// RUN: %clang_cc1 -fsycl-is-host %s -verify -DUNNAMED_LAMBDAS // // RUN: %clang_cc1 -fsycl-is-host -fno-sycl-unnamed-lambda %s -verify -// RUN: %clang_cc1 -fsycl-is-host -sycl-std=2017 -fno-sycl-unnamed-lambda %s -verify -// RUN: %clang_cc1 -fsycl-is-host -sycl-std=2017 -fsycl-unnamed-lambda %s -verify -DUNNAMED_LAMBDAS +// RUN: %clang_cc1 -fsycl-is-host -fno-sycl-unnamed-lambda %s -verify +// RUN: %clang_cc1 -fsycl-is-host -fsycl-unnamed-lambda %s -verify -DUNNAMED_LAMBDAS // expected-no-diagnostics #if defined(UNNAMED_LAMBDAS) && !defined(__SYCL_UNNAMED_LAMBDA__) -#error "Unnamed lambdas should be enabled for SYCL2020/default" +#error "Unnamed lambdas should be enabled" #endif #if !defined(UNNAMED_LAMBDAS) && defined(__SYCL_UNNAMED_LAMBDA__) -#error "Unnamed lambdas should NOT be enabled for SYCL2017, or when explicitly disabled" +#error "Unnamed lambdas should NOT be enabled" #endif diff --git a/clang/test/Preprocessor/sycl-macro.cpp b/clang/test/Preprocessor/sycl-macro.cpp index 81645d3ac373..f97ba32db86e 100644 --- a/clang/test/Preprocessor/sycl-macro.cpp +++ b/clang/test/Preprocessor/sycl-macro.cpp @@ -1,10 +1,7 @@ // RUN: %clang_cc1 %s -E -dM | FileCheck %s // RUN: %clang_cc1 %s -fsycl-is-device -E -dM | FileCheck --check-prefix=CHECK-SYCL-ID %s -// RUN: %clang_cc1 %s -fsycl-id-queries-fit-in-int -fsycl-is-host -sycl-std=2017 -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD %s -// RUN: %clang_cc1 %s -fsycl-id-queries-fit-in-int -fsycl-is-device -sycl-std=2017 -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD %s -// RUN: %clang_cc1 %s -fsycl-id-queries-fit-in-int -fsycl-is-host -sycl-std=2020 -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD-2020 %s // RUN: %clang_cc1 %s -fsycl-id-queries-fit-in-int -fsycl-is-host -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD-2020 %s -// RUN: %clang_cc1 %s -fsycl-id-queries-fit-in-int -fsycl-is-device -sycl-std=2020 -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD-2020 %s +// RUN: %clang_cc1 %s -fsycl-id-queries-fit-in-int -fsycl-is-device -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD-2020 %s // RUN: %clang_cc1 %s -fsycl-id-queries-fit-in-int -fsycl-is-device -sycl-std=1.2.1 -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD-DEVICE %s // RUNx: %clang_cc1 %s -fsycl-id-queries-fit-in-int -fsycl-is-device -E -dM -fms-compatibility | FileCheck --check-prefix=CHECK-MSVC %s // RUN: %clang_cc1 -fno-sycl-id-queries-fit-in-int %s -E -dM | FileCheck \ @@ -19,11 +16,12 @@ // CHECK-NOT:#define CL_SYCL_LANGUAGE_VERSION 121 // CHECK-NOT:#define __SYCL_ID_QUERIES_FIT_IN_INT__ 1 -// CHECK-SYCL-STD:#define CL_SYCL_LANGUAGE_VERSION 121 -// CHECK-SYCL-STD:#define SYCL_LANGUAGE_VERSION 201707 +// CHECK-SYCL-STD:#define SYCL_LANGUAGE_VERSION 202001 + // CHECK-SYCL-STD:#define __SYCL_ID_QUERIES_FIT_IN_INT__ 1 // CHECK-SYCL-STD-2020:#define SYCL_LANGUAGE_VERSION 202001 +// CHECK-SYCL-STD-2020:#define __SYCL_ID_QUERIES_FIT_IN_INT__ 1 // CHECK-SYCL-STD-DEVICE:#define __SYCL_DEVICE_ONLY__ 1 // CHECK-SYCL-STD-DEVICE:#define __SYCL_ID_QUERIES_FIT_IN_INT__ 1 diff --git a/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp b/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp deleted file mode 100644 index f40bf981ea60..000000000000 --- a/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp +++ /dev/null @@ -1,293 +0,0 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump %s | FileCheck %s - -// Tests for AST of [[intel::scheduler_target_fmax_mhz()]], [[intel::num_simd_work_items()]], -// [[intel::no_global_work_offset()]], [[intel::max_global_work_dim()]], [[intel::sycl_explicit_simd]], -// [[sycl::reqd_sub_group_size()]], [[sycl::reqd_work_group_size()]], [[intel::kernel_args_restrict]], and -// [[intel::max_work_group_size()]] function attributes in SYCL 2020. - -#include "sycl.hpp" - -sycl::queue deviceQueue; - -struct FuncObj { - [[intel::sycl_explicit_simd]] void operator()() const {} -}; - -struct FuncObj1 { - [[intel::no_global_work_offset(1)]] void operator()() const {} -}; - -struct FuncObj2 { - [[intel::scheduler_target_fmax_mhz(10)]] void operator()() const {} -}; - -struct FuncObj3 { - [[intel::max_work_group_size(2, 2, 2)]] void operator()() const {} -}; - -struct FuncObj4 { - [[sycl::reqd_work_group_size(2, 2, 2)]] void operator()() const {} -}; - -struct FuncObj5 { - [[intel::num_simd_work_items(8)]] void operator()() const {} -}; - -struct FuncObj6 { - [[intel::kernel_args_restrict]] void operator()() const {} -}; - -struct FuncObj7 { - [[intel::max_global_work_dim(1)]] void operator()() const {} -}; - -[[intel::sycl_explicit_simd]] void func() {} - -[[intel::no_global_work_offset(1)]] void func1() {} - -[[intel::scheduler_target_fmax_mhz(2)]] void func2() {} - -[[intel::max_work_group_size(1, 1, 1)]] void func3() {} - -[[sycl::reqd_work_group_size(1, 1, 1)]] void func4() {} - -[[intel::num_simd_work_items(5)]] void func5() {} - -[[intel::kernel_args_restrict]] void func6() {} - -[[intel::max_global_work_dim(0)]] void func7() {} - -[[sycl::reqd_sub_group_size(4)]] void func8() {} - -class Functor { -public: - void operator()() const { - func8(); - } -}; - -class Functor1 { -public: - [[sycl::reqd_sub_group_size(12)]] void operator()() const {} -}; - -int main() { - deviceQueue.submit([&](sycl::handler &h) { - // CHECK: FunctionDecl {{.*}}test_kernel1 - // CHECK: SYCLSimdAttr {{.*}} Implicit - // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit - // CHECK-NEXT: AsmLabelAttr {{.*}} Implicit - // CHECK-NEXT: SYCLSimdAttr - h.single_task( - FuncObj()); - // CHECK: FunctionDecl {{.*}}test_kernel2 - // CHECK: SYCLSimdAttr {{.*}} Implicit - // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit - // CHECK-NEXT: AsmLabelAttr {{.*}} Implicit - // CHECK-NEXT: SYCLSimdAttr - h.single_task( - []() [[intel::sycl_explicit_simd]]{}); - - // Test attribute is not propagated. - // CHECK: FunctionDecl {{.*}}test_kernel3 - // CHECK: SYCLSimdAttr {{.*}} Implicit - // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit - // CHECK-NEXT: AsmLabelAttr {{.*}} Implicit - // CHECK-NEXT: SYCLSimdAttr - // CHECK-NOT: SYCLSimdAttr - h.single_task( - []() [[intel::sycl_explicit_simd]] { func(); }); - - // Test attribute is not propagated. - // CHECK: FunctionDecl {{.*}}test_kernel4 - // CHECK-NOT: SYCLIntelNoGlobalWorkOffsetAttr - h.single_task( - []() { func1(); }); - - // CHECK: FunctionDecl {{.*}}test_kernel5 - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 1 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 - h.single_task( - FuncObj1()); - - // CHECK: FunctionDecl {{.*}}test_kernel6 - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 1 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 - h.single_task( - []() [[intel::no_global_work_offset]]{}); - - // CHECK: FunctionDecl {{.*}}test_kernel7 - // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 10 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 10 - h.single_task( - FuncObj2()); - - // CHECK: FunctionDecl {{.*}}test_kernel8 - // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 20 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 20 - h.single_task( - []() [[intel::scheduler_target_fmax_mhz(20)]]{}); - - // Test attribute is not propagated. - // CHECK: FunctionDecl {{.*}}test_kernel9 - // CHECK-NOT: SYCLIntelSchedulerTargetFmaxMhzAttr - h.single_task( - []() { func2(); }); - - // CHECK: FunctionDecl {{.*}}test_kernel10 - // CHECK: SYCLIntelMaxWorkGroupSizeAttr - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 - h.single_task( - FuncObj3()); - - // Test attribute is not propagated. - // CHECK: FunctionDecl {{.*}}test_kernel11 - // CHECK-NOT: SYCLIntelMaxWorkGroupSizeAttr - h.single_task( - []() { func3(); }); - - // CHECK: FunctionDecl {{.*}}test_kernel12 - // CHECK: SYCLIntelMaxWorkGroupSizeAttr - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 - h.single_task( - []() [[intel::max_work_group_size(8, 8, 8)]]{}); - - // CHECK: FunctionDecl {{.*}}test_kernel13 - // CHECK: SYCLReqdWorkGroupSizeAttr - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 - h.single_task( - FuncObj4()); - - // Test attribute is not propagated. - // CHECK: FunctionDecl {{.*}}test_kernel14 - // CHECK-NOT: SYCLReqdWorkGroupSizeAttr - h.single_task( - []() { func4(); }); - - // CHECK: FunctionDecl {{.*}}test_kernel15 - // CHECK: SYCLReqdWorkGroupSizeAttr - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 - h.single_task( - []() [[sycl::reqd_work_group_size(8, 8, 8)]]{}); - - // CHECK: FunctionDecl {{.*}}test_kernel16 - // CHECK: SYCLIntelNumSimdWorkItemsAttr - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 - h.single_task( - FuncObj5()); - - // CHECK: FunctionDecl {{.*}}test_kernel17 - // CHECK: SYCLIntelNumSimdWorkItemsAttr - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 20 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 20 - h.single_task( - []() [[intel::num_simd_work_items(20)]]{}); - - // Test attribute is not propagated. - // CHECK: FunctionDecl {{.*}}test_kernel18 - // CHECK-NOT: SYCLIntelNumSimdWorkItemsAttr - h.single_task( - []() { func5(); }); - - // CHECK: FunctionDecl {{.*}}test_kernel19 - // CHECK: SYCLIntelKernelArgsRestrictAttr - h.single_task( - FuncObj6()); - - // CHECK: FunctionDecl {{.*}}test_kernel20 - // CHECK: SYCLIntelKernelArgsRestrictAttr - h.single_task( - []() [[intel::kernel_args_restrict]]{}); - - // Test attribute is not propagated. - // CHECK: FunctionDecl {{.*}}test_kernel21 - // CHECK-NOT: SYCLIntelKernelArgsRestrictAttr - h.single_task( - []() { func6(); }); - - // CHECK: FunctionDecl {{.*}}test_kernel22 - // CHECK: SYCLIntelMaxGlobalWorkDimAttr - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 1 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 - h.single_task( - FuncObj7()); - - // CHECK: FunctionDecl {{.*}}test_kernel23 - // CHECK: SYCLIntelMaxGlobalWorkDimAttr - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 0 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 0 - h.single_task( - []() [[intel::max_global_work_dim(0)]]{}); - - // Test attribute is not propagated. - // CHECK: FunctionDecl {{.*}}test_kernel24 - // CHECK-NOT: SYCLIntelMaxGlobalWorkDimAttr - h.single_task( - []() { func7(); }); - - // Test attribute is not propagated. - // CHECK: FunctionDecl {{.*}}test_kernel25 - // CHECK-NOT: IntelReqdSubGroupSizeAttr - Functor f; - h.single_task(f); - - // CHECK: FunctionDecl {{.*}}test_kernel26 - // CHECK: IntelReqdSubGroupSizeAttr - Functor1 f1; - h.single_task(f1); - - // CHECK: FunctionDecl {{.*}}test_kernel27 - // CHECK: IntelReqdSubGroupSizeAttr - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 - h.single_task( - []() [[sycl::reqd_sub_group_size(8)]]{}); - }); - return 0; -} diff --git a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp b/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp deleted file mode 100644 index bdabd1f40d52..000000000000 --- a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp +++ /dev/null @@ -1,54 +0,0 @@ -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 -verify -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 -DTRIGGER_ERROR -verify -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -sycl-std=2017 -triple spir64 | FileCheck %s - -#ifndef TRIGGER_ERROR -[[intel::no_global_work_offset]] void not_direct_one() {} // expected-no-diagnostics - -[[intel::reqd_sub_group_size(1)]] void func_one() { - not_direct_one(); -} - -#else -[[sycl::reqd_work_group_size(2, 2, 2)]] void not_direct_two() {} // expected-note 2 {{conflicting attribute is here}} - -[[intel::max_work_group_size(1, 1, 1)]] // expected-note 3 {{conflicting attribute is here}} -void -func_two() { - not_direct_two(); -} - -[[sycl::reqd_work_group_size(4, 4, 4)]] // expected-note 1 {{conflicting attribute is here}} -void -func_three() { - not_direct_two(); -} -#endif - -template -[[clang::sycl_kernel]] void __my_kernel__(const Type &bar) { - bar(); -#ifndef TRIGGER_ERROR - func_one(); -#else - func_two(); - func_three(); -#endif -} - -template -void parallel_for(Type lambda) { - __my_kernel__(lambda); -} - -void invoke_foo2() { -#ifndef TRIGGER_ERROR - // CHECK-LABEL: FunctionDecl {{.*}} invoke_foo2 'void ()' - // CHECK: `-FunctionDecl {{.*}}KernelName 'void ()' - // CHECK: -IntelReqdSubGroupSizeAttr {{.*}} - // CHECK: `-SYCLIntelNoGlobalWorkOffsetAttr {{.*}} - parallel_for([]() {}); -#else - parallel_for([]() {}); // expected-error 3 {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}} -#endif -} diff --git a/clang/test/SemaSYCL/check-work-group-size-hint-device.cpp b/clang/test/SemaSYCL/check-work-group-size-hint-device.cpp index b989bb3d8bb2..a0cbcef3e83b 100644 --- a/clang/test/SemaSYCL/check-work-group-size-hint-device.cpp +++ b/clang/test/SemaSYCL/check-work-group-size-hint-device.cpp @@ -1,10 +1,10 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify -DEXPECT_PROP -DTRIGGER_ERROR %s -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -ast-dump -DEXPECT_PROP %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -verify -DEXPECT_PROP -DTRIGGER_ERROR %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump -DEXPECT_PROP %s | FileCheck %s -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2020 -verify -DTRIGGER_ERROR %s -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -verify -DTRIGGER_ERROR %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump %s | FileCheck %s -// Test for AST of work_group_size_hint kernel attribute in SYCL 1.2.1. and SYCL 2020 modes. +// Test for AST of work_group_size_hint kernel attribute in SYCL. #include "sycl.hpp" // Check the basics. @@ -133,28 +133,6 @@ void invoke() { // CHECK-NEXT: value: Int 1 // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} - // Checking that attributes are propagated to the kernel from functions in SYCL 1.2.1 mode. -#ifdef EXPECT_PROP - h.single_task([=]() { - f4x4x4(); - }); -#else - // Otherwise using a functor that has the required attributes - h.single_task(f4x4x4); -#endif - // CHECK: FunctionDecl {{.*}} {{.*}}kernel_2 - // CHECK: SYCLWorkGroupSizeHintAttr {{.*}} - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 4 - // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 4 - // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 4 - // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} - - // Check that conflicts are reported if the attribute is propagated in SYCL 1.2.1 mode. FunctorNoProp fNoProp; h.single_task(fNoProp); diff --git a/clang/test/SemaSYCL/implicit_kernel_type.cpp b/clang/test/SemaSYCL/implicit_kernel_type.cpp index 3343c9c6d7fd..84f8ac12574f 100644 --- a/clang/test/SemaSYCL/implicit_kernel_type.cpp +++ b/clang/test/SemaSYCL/implicit_kernel_type.cpp @@ -1,6 +1,6 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fno-sycl-unnamed-lambda -fsyntax-only -sycl-std=2020 -verify %s -Werror=sycl-strict -DERROR -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fno-sycl-unnamed-lambda -fsyntax-only -sycl-std=2020 -verify %s -Wsycl-strict -DWARN -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2020 -verify %s -Werror=sycl-strict -DERROR +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fno-sycl-unnamed-lambda -fsyntax-only -verify %s -Werror=sycl-strict -DERROR +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fno-sycl-unnamed-lambda -fsyntax-only -verify %s -Wsycl-strict -DWARN +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -verify %s -Werror=sycl-strict -DERROR // This test verifies that incorrect kernel names are diagnosed correctly. diff --git a/clang/test/SemaSYCL/intel-max-global-work-dim-device-ast.cpp b/clang/test/SemaSYCL/intel-max-global-work-dim-device-ast.cpp index 2de88108651b..b8a171c83289 100644 --- a/clang/test/SemaSYCL/intel-max-global-work-dim-device-ast.cpp +++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device-ast.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple spir64 | FileCheck %s +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 | FileCheck %s // The test checks AST of [[intel::max_global_work_dim()]] attribute. @@ -55,8 +55,6 @@ int ver() { return 0; } -[[intel::max_global_work_dim(2)]] void func_do_not_ignore() {} - struct FuncObj { [[intel::max_global_work_dim(1)]] void operator()() const {} }; @@ -133,14 +131,6 @@ int main() { h.single_task( []() [[intel::max_global_work_dim(2)]] {}); - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 - // CHECK: SYCLIntelMaxGlobalWorkDimAttr - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} - h.single_task( - []() { func_do_not_ignore(); }); - h.single_task(TRIFuncObjGood1()); // CHECK-LABEL: FunctionDecl {{.*}}test_kernel4 // CHECK: SYCLIntelMaxGlobalWorkDimAttr diff --git a/clang/test/SemaSYCL/intel-max-work-group-size-ast.cpp b/clang/test/SemaSYCL/intel-max-work-group-size-ast.cpp index 59f80bb24695..0c1aa241851b 100644 --- a/clang/test/SemaSYCL/intel-max-work-group-size-ast.cpp +++ b/clang/test/SemaSYCL/intel-max-work-group-size-ast.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple spir64 | FileCheck %s +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 | FileCheck %s // The test checks support and functionality of [[intel:::max_work_group_size()]] attribute. #include "sycl.hpp" @@ -6,8 +6,6 @@ using namespace sycl; queue q; -[[intel::max_work_group_size(2, 2, 2)]] void func_do_not_ignore() {} - struct FuncObj { [[intel::max_work_group_size(4, 4, 4)]] void operator()() const {} }; @@ -96,20 +94,6 @@ int main() { h.single_task( []() [[intel::max_work_group_size(8, 8, 8)]] {}); - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 - // CHECK: SYCLIntelMaxWorkGroupSizeAttr - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} - h.single_task( - []() { func_do_not_ignore(); }); - // Ignore duplicate attribute. h.single_task( // CHECK-LABEL: FunctionDecl {{.*}}test_kernel10 diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size-ast-device.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size-ast-device.cpp index 5ce303af1cb5..88c767f1175a 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size-ast-device.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size-ast-device.cpp @@ -1,6 +1,6 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump %s | FileCheck %s -// Test for AST of reqd_work_group_size kernel attribute in SYCL 1.2.1. +// Test for AST of reqd_work_group_size kernel attribute in SYCL. #include "sycl.hpp" @@ -65,8 +65,6 @@ int check() { // CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} -[[sycl::reqd_work_group_size(4)]] void f4() {} - class Functor16 { public: [[sycl::reqd_work_group_size(16)]] void operator()() const {} @@ -77,25 +75,11 @@ class Functor16x16x16 { [[sycl::reqd_work_group_size(16, 16, 16)]] void operator()() const {} }; -class Functor { -public: - void operator()() const { - f4(); - } -}; - class FunctorAttr { public: [[sycl::reqd_work_group_size(128, 128, 128)]] void operator()() const {} }; -// Test of redeclaration of [[intel::max_work_group_size()]] and [[sycl::reqd_work_group_size()]]. -[[intel::no_global_work_offset]] void func1(); -[[intel::max_work_group_size(4, 4, 4)]] void func1(); -[[sycl::reqd_work_group_size(2, 2, 2)]] void func1() {} - -[[sycl::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} - int main() { q.submit([&](handler &h) { // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 @@ -106,14 +90,6 @@ int main() { Functor16 f16; h.single_task(f16); - // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2 - // CHECK: SYCLReqdWorkGroupSizeAttr - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 4 - // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} - Functor f; - h.single_task(f); - // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name3 // CHECK: SYCLReqdWorkGroupSizeAttr // CHECK-NEXT: ConstantExpr{{.*}}'int' @@ -142,7 +118,7 @@ int main() { FunctorAttr fattr; h.single_task(fattr); - // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 +// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 // CHECK: SYCLReqdWorkGroupSizeAttr // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 32 @@ -153,37 +129,7 @@ int main() { // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 32 // CHECK-NEXT: IntegerLiteral{{.*}}32{{$}} - h.single_task([]() [[sycl::reqd_work_group_size(32, 32, 32)]] { - f32x32x32(); - }); - - // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name6 - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 1 - // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} - // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 4 - // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 4 - // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 4 - // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} - // CHECK: SYCLReqdWorkGroupSizeAttr - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} - h.single_task( - []() { func1(); }); + h.single_task([]() [[sycl::reqd_work_group_size(32, 32, 32)]] {}); }); return 0; } diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp deleted file mode 100644 index 88ac8380a69b..000000000000 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp +++ /dev/null @@ -1,67 +0,0 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -fsyntax-only -verify %s - -// The test checks support and functionality of reqd_work_group_size kernel attribute in SYCL 2017. - -#include "sycl.hpp" - -using namespace sycl; -queue q; - -[[sycl::reqd_work_group_size(4)]] void f4() {} // expected-note {{conflicting attribute is here}} -// expected-note@-1 {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(32)]] void f32() {} // expected-note {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(16)]] void f16() {} // expected-note {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(16, 16)]] void f16x16() {} // expected-note {{conflicting attribute is here}} - -[[sycl::reqd_work_group_size(32, 32)]] void f32x32() {} // expected-note {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} - -[[intel::reqd_work_group_size(4, 2, 9)]] void unknown() {} // expected-warning{{unknown attribute 'reqd_work_group_size' ignored}} - -class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} -public: - [[sycl::reqd_work_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}} - f4(); - } -}; - -// Tests of redeclaration of [[intel::max_work_group_size()]] and [[sycl::reqd_work_group_size()]] - expect error -[[intel::max_work_group_size(4, 4, 4)]] void func2(); // expected-note {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(8, 8, 8)]] void func2() {} // expected-error {{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} - -[[sycl::reqd_work_group_size(4, 4, 4)]] void func3(); // expected-note {{previous attribute is here}} -[[sycl::reqd_work_group_size(1, 1, 1)]] void func3() {} // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} - -int main() { - q.submit([&](handler &h) { - Functor8 f8; - h.single_task(f8); - - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} - f4(); - f32(); - }); - - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} - f16(); - f16x16(); - }); - - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} - f32x32x32(); - f32x32(); - }); - - // expected-error@+1 {{expected variable name or 'this' in lambda capture list}} - h.single_task([[sycl::reqd_work_group_size(32, 32, 32)]][]() { - f32x32x32(); - }); - - h.single_task( - []() { func2(); }); - - h.single_task( - []() { func3(); }); - }); - return 0; -} diff --git a/clang/test/SemaSYCL/intel-restrict.cpp b/clang/test/SemaSYCL/intel-restrict.cpp index f8e7a670fddd..1ae5ff1490b3 100644 --- a/clang/test/SemaSYCL/intel-restrict.cpp +++ b/clang/test/SemaSYCL/intel-restrict.cpp @@ -1,7 +1,5 @@ -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 -DCHECKDIAG -verify -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 | FileCheck %s - -[[intel::kernel_args_restrict]] void func_do_not_ignore() {} +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -triple spir64 -DCHECKDIAG -verify +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -triple spir64 | FileCheck %s struct FuncObj { [[intel::kernel_args_restrict]] void operator()() const {} @@ -26,8 +24,4 @@ int main() { kernel( []() [[intel::kernel_args_restrict]] {}); - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 - // CHECK: SYCLIntelKernelArgsRestrictAttr - kernel( - []() { func_do_not_ignore(); }); } diff --git a/clang/test/SemaSYCL/intel-work-group-size-hint-ast-device.cpp b/clang/test/SemaSYCL/intel-work-group-size-hint-ast-device.cpp index caf429e5a7bc..7cfd108f6ad8 100644 --- a/clang/test/SemaSYCL/intel-work-group-size-hint-ast-device.cpp +++ b/clang/test/SemaSYCL/intel-work-group-size-hint-ast-device.cpp @@ -1,6 +1,6 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump %s | FileCheck %s -// Test for AST of work_group_size_hint kernel attribute in SYCL 1.2.1. +// Test for AST of work_group_size_hint kernel attribute in SYCL. #include "sycl.hpp" @@ -65,8 +65,6 @@ int check() { // CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} -[[sycl::work_group_size_hint(4)]] void f4() {} - class Functor16 { public: [[sycl::work_group_size_hint(16)]] void operator()() const {} @@ -77,25 +75,11 @@ class Functor32x16x8 { [[sycl::work_group_size_hint(32, 16, 8)]] void operator()() const {} }; -class Functor { -public: - void operator()() const { - f4(); - } -}; - class FunctorAttr { public: [[sycl::work_group_size_hint(128, 256, 512)]] void operator()() const {} }; -// Test of redeclaration of [[intel::max_work_group_size()]] and [[sycl::work_group_size_hint()]]. -[[intel::no_global_work_offset]] void func1(); -[[intel::max_work_group_size(4, 2, 6)]] void func1(); -[[sycl::work_group_size_hint(2, 1, 3)]] void func1() {} - -[[sycl::work_group_size_hint(8, 16, 32)]] void f8x16x32() {} - int main() { q.submit([&](handler &h) { // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 @@ -106,14 +90,6 @@ int main() { Functor16 f16; h.single_task(f16); - // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2 - // CHECK: SYCLWorkGroupSizeHintAttr - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 4 - // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} - Functor f; - h.single_task(f); - // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name3 // CHECK: SYCLWorkGroupSizeHintAttr // CHECK-NEXT: ConstantExpr{{.*}}'int' @@ -142,48 +118,6 @@ int main() { FunctorAttr fattr; h.single_task(fattr); - // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 - // CHECK: SYCLWorkGroupSizeHintAttr - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 16 - // CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 32 - // CHECK-NEXT: IntegerLiteral{{.*}}32{{$}} - h.single_task([]() [[sycl::work_group_size_hint(8, 16, 32)]] { - f8x16x32(); - }); - - // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name6 - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 1 - // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} - // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 4 - // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 6 - // CHECK-NEXT: IntegerLiteral{{.*}}6{{$}} - // CHECK: SYCLWorkGroupSizeHintAttr - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 1 - // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 3 - // CHECK-NEXT: IntegerLiteral{{.*}}3{{$}} - h.single_task( - []() { func1(); }); }); return 0; } diff --git a/clang/test/SemaSYCL/lb_sm_90_ast.cpp b/clang/test/SemaSYCL/lb_sm_90_ast.cpp index 6a53adeb6c99..fb5c0f383abc 100644 --- a/clang/test/SemaSYCL/lb_sm_90_ast.cpp +++ b/clang/test/SemaSYCL/lb_sm_90_ast.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -ast-dump -triple nvptx-unknown-unknown -target-cpu sm_90 -Wno-c++23-extensions %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump -triple nvptx-unknown-unknown -target-cpu sm_90 -Wno-c++23-extensions %s | FileCheck %s // Tests for AST of Intel max_work_group_size, min_work_groups_per_cu and // max_work_groups_per_mp attribute. @@ -7,30 +7,6 @@ sycl::queue deviceQueue; -// CHECK: FunctionDecl {{.*}} func1 'void ()' -// CHECK-NEXT: CompoundStmt -// CHECK-NEXT: SYCLIntelMaxWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 8 -// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 8 -// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 8 -// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 -// CHECK-NEXT: SYCLIntelMinWorkGroupsPerComputeUnitAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 4 -// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 4 -// CHECK-NEXT: SYCLIntelMaxWorkGroupsPerMultiprocessorAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 2 -// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 -[[intel::max_work_group_size(8, 8, 8), intel::min_work_groups_per_cu(4), - intel::max_work_groups_per_mp(2)]] void -func1() {} - // Test that checks template parameter support on function. // CHECK: FunctionTemplateDecl {{.*}} func2 // CHECK: FunctionDecl {{.*}} func2 'void ()' @@ -80,11 +56,6 @@ template intel::max_work_groups_per_mp(N)]] void func2() {} -class KernelFunctor { -public: - void operator()() const { func1(); } -}; - // Test that checks template parameter support on class member function. template class KernelFunctor2 { public: @@ -95,27 +66,6 @@ template class KernelFunctor2 { int main() { deviceQueue.submit([&](sycl::handler &h) { - // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_1 - // CHECK: SYCLIntelMaxWorkGroupSizeAttr - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 - // CHECK: SYCLIntelMinWorkGroupsPerComputeUnitAttr - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 4 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 4 - // CHECK: SYCLIntelMaxWorkGroupsPerMultiprocessorAttr - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 - KernelFunctor f1; - h.single_task(f1); // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_2 // CHECK: SYCLIntelMaxWorkGroupSizeAttr diff --git a/clang/test/SemaSYCL/num_simd_work_items_ast.cpp b/clang/test/SemaSYCL/num_simd_work_items_ast.cpp index c2dbc246c651..4c66a6380cd1 100644 --- a/clang/test/SemaSYCL/num_simd_work_items_ast.cpp +++ b/clang/test/SemaSYCL/num_simd_work_items_ast.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 %s -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -fsyntax-only -sycl-std=2017 -ast-dump | FileCheck %s +// RUN: %clang_cc1 %s -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -fsyntax-only -ast-dump | FileCheck %s // The test checks AST of [[intel::num_simd_work_items()]] attribute. @@ -54,8 +54,6 @@ int ver() { return 0; } -[[intel::num_simd_work_items(2)]] void func_do_not_ignore() {} - struct FuncObj { [[intel::num_simd_work_items(42)]] void operator()() const {} }; @@ -116,14 +114,6 @@ int main() { h.single_task( []() [[intel::num_simd_work_items(8)]] {}); - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 - // CHECK: SYCLIntelNumSimdWorkItemsAttr {{.*}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} - h.single_task( - []() { func_do_not_ignore(); }); - h.single_task(TRIFuncObjGood1()); // CHECK-LABEL: FunctionDecl {{.*}}test_kernel4 // CHECK: SYCLIntelNumSimdWorkItemsAttr {{.*}} diff --git a/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp b/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp index ee86f1b124c2..c4da05dfbf23 100755 --- a/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp +++ b/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -sycl-std=2017 -fsycl-is-device -triple spir64 | FileCheck %s +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -triple spir64 | FileCheck %s #include "Inputs/sycl.hpp" diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp index afb45749c626..8386663bbd37 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump %s | FileCheck %s // The test checks AST of [[intel::reqd_sub_group_size()]] attribute. @@ -7,20 +7,11 @@ using namespace sycl; queue q; -[[intel::reqd_sub_group_size(4)]] void foo() {} - class Functor16 { public: [[intel::reqd_sub_group_size(16)]] void operator()() const {} }; -class Functor { -public: - void operator()() const { - foo(); - } -}; - // Test that checks template parameter support on member function of class template. template class KernelFunctor { @@ -56,14 +47,6 @@ int main() { Functor16 f16; h.single_task(f16); - // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2 - // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 4 - // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} - Functor f; - h.single_task(f); - // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name3 // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size // CHECK-NEXT: ConstantExpr {{.*}} 'int' @@ -71,12 +54,6 @@ int main() { // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} h.single_task([]() [[intel::reqd_sub_group_size(2)]] {}); - // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name4 - // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 4 - // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} - h.single_task([]() [[intel::reqd_sub_group_size(4)]] { foo(); }); // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size // CHECK-NEXT: ConstantExpr {{.*}} 'int' diff --git a/clang/test/SemaSYCL/reqd-sub-group-size.cpp b/clang/test/SemaSYCL/reqd-sub-group-size.cpp index 85afcf97388f..218527972128 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size.cpp @@ -1,32 +1,34 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify -pedantic %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -verify -pedantic %s // The test checks functionality of [[intel::reqd_sub_group_size()]] attribute on SYCL kernel. -#include "sycl.hpp" -using namespace sycl; -queue q; +#include "sycl.hpp" //clang/test/SemaSYCL/Inputs/sycl.hpp -[[intel::reqd_sub_group_size(4)]] void foo() {} // expected-note {{conflicting attribute is here}} -// expected-note@-1 {{conflicting attribute is here}} -[[intel::reqd_sub_group_size(32)]] void baz() {} // expected-note {{conflicting attribute is here}} +sycl::queue q; -class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} -public: - [[intel::reqd_sub_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}} - foo(); - } +// Kernel defined as a named function object +class KernelFunctor1 { + public: + [[sycl::reqd_work_group_size(16)]] void operator()() const {}; +}; + +// Kernel defined as a named function object +class KernelFunctor2 { + public: + void operator() [[sycl::reqd_work_group_size(16)]] () const {}; }; int main() { - q.submit([&](handler &h) { - Functor8 f8; - h.single_task(f8); - - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} - foo(); - baz(); - }); + // Kernel defined as a lambda + q.submit([&](sycl::handler& h) { + KernelFunctor1 kf1; + KernelFunctor2 kf2; + h.single_task(kf1); + h.single_task(kf2); + h.single_task( + []()[[sycl::reqd_work_group_size(16)]]{} + ); }); return 0; } diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp deleted file mode 100644 index f65a1b59909c..000000000000 --- a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp +++ /dev/null @@ -1,186 +0,0 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify -DTRIGGER_ERROR %s -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s - -// Test for AST of reqd_work_group_size kernel attribute in SYCL 1.2.1. -#include "sycl.hpp" - -using namespace sycl; -queue q; - -[[sycl::reqd_work_group_size(4, 1, 1)]] void f4x1x1() {} // expected-note {{conflicting attribute is here}} -// expected-note@-1 {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(32, 1, 1)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}} - -[[sycl::reqd_work_group_size(16, 1, 1)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(16, 16, 1)]] void f16x16x1() {} // expected-note {{conflicting attribute is here}} - -[[sycl::reqd_work_group_size(32, 32, 1)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} - -// No diagnostic because the attributes are synonyms with identical behavior. -[[sycl::reqd_work_group_size(4, 4, 4)]] void four1(); -[[sycl::reqd_work_group_size(4, 4, 4)]] void four1(); // OK -[[sycl::reqd_work_group_size(4, 4)]] void four2(); -[[sycl::reqd_work_group_size(4, 4)]] void four2(); // OK -[[sycl::reqd_work_group_size(4)]] void four3(); -[[sycl::reqd_work_group_size(4)]] void four3(); // OK - -#ifdef TRIGGER_ERROR -// Althrough optional values are effectively representing 1's, the attributes -// should be considered different when there is a different in the arguments -// given. -[[sycl::reqd_work_group_size(4)]] void four4(); // expected-note {{previous attribute is here}} -[[sycl::reqd_work_group_size(4, 1)]] void four4(); // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} -[[sycl::reqd_work_group_size(4)]] void four5(); // expected-note {{previous attribute is here}} -[[sycl::reqd_work_group_size(4, 1, 1)]] void four5(); // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} -[[sycl::reqd_work_group_size(4)]] void four6(); // expected-note {{previous attribute is here}} -[[sycl::reqd_work_group_size(1, 4)]] void four6(); // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} -[[sycl::reqd_work_group_size(4)]] void four7(); // expected-note {{previous attribute is here}} -[[sycl::reqd_work_group_size(1, 1, 4)]] void four7(); // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} -#endif - -// Make sure there's at least one argument passed for the SYCL spelling. -#ifdef TRIGGER_ERROR -[[sycl::reqd_work_group_size]] void four_no_more(); // expected-error {{'reqd_work_group_size' attribute takes at least 1 argument}} -#endif // TRIGGER_ERROR - -class Functor16 { -public: - [[sycl::reqd_work_group_size(16, 1, 1)]] [[sycl::reqd_work_group_size(16, 1, 1)]] void operator()() const {} -}; - -#ifdef TRIGGER_ERROR -class Functor32 { -public: - [[sycl::reqd_work_group_size(32, 1, 1)]] // expected-note {{previous attribute is here}} - [[sycl::reqd_work_group_size(1, 1, 32)]] void // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} - operator()() const {} -}; -#endif -class Functor16x16x16 { -public: - [[sycl::reqd_work_group_size(16, 16, 16)]] void operator()() const {} -}; - -class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} -public: - [[sycl::reqd_work_group_size(1, 1, 8)]] void operator()() const { // expected-note {{conflicting attribute is here}} - f4x1x1(); - } -}; - -class Functor { -public: - void operator()() const { - f4x1x1(); - } -}; - -int main() { - q.submit([&](handler &h) { - Functor16 f16; - h.single_task(f16); - - Functor f; - h.single_task(f); - - Functor16x16x16 f16x16x16; - h.single_task(f16x16x16); - - h.single_task([]() [[sycl::reqd_work_group_size(32, 32, 32), sycl::reqd_work_group_size(32, 32, 32)]] { - f32x32x32(); - }); - -#ifdef TRIGGER_ERROR - Functor8 f8; - h.single_task(f8); - - Functor32 f32; - h.single_task(f32); - - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} - f4x1x1(); - f32x1x1(); - }); - - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} - f16x1x1(); - f16x16x1(); - }); - - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} - f32x32x32(); - f32x32x1(); - }); - - // expected-error@+1 {{expected variable name or 'this' in lambda capture list}} - h.single_task([[sycl::reqd_work_group_size(32, 32, 32)]][]() { - f32x32x32(); - }); - -#endif - // Ignore duplicate attribute. - h.single_task( - []() [[sycl::reqd_work_group_size(2, 2, 2), - sycl::reqd_work_group_size(2, 2, 2)]] {}); - }); - return 0; -} - -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 -// CHECK: SYCLReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 16 -// CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 1 -// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 1 -// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2 -// CHECK: SYCLReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 4 -// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 1 -// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 1 -// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name3 -// CHECK: SYCLReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 16 -// CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 16 -// CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 16 -// CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 -// CHECK: SYCLReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 32 -// CHECK-NEXT: IntegerLiteral{{.*}}32{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 32 -// CHECK-NEXT: IntegerLiteral{{.*}}32{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 32 -// CHECK-NEXT: IntegerLiteral{{.*}}32{{$}} -// -// CHECK: FunctionDecl {{.*}}test_kernel11 -// CHECK: SYCLReqdWorkGroupSizeAttr -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 2 -// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 2 -// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 2 -// CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} -// CHECK-NOT: SYCLReqdWorkGroupSizeAttr diff --git a/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp b/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp index 075cd9877d04..09aa899e7fbf 100644 --- a/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp +++ b/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp @@ -1,18 +1,10 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump %s | FileCheck %s // Tests for AST of Intel FPGA scheduler_target_fmax_mhz function attribute. #include "sycl.hpp" sycl::queue deviceQueue; -// CHECK: FunctionDecl {{.*}} func1 'void ()' -// CHECK-NEXT: CompoundStmt -// CHECK-NEXT: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 4 -// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 4 -[[intel::scheduler_target_fmax_mhz(4)]] void func1() {} - // Test that checks template parameter support on function. // CHECK: FunctionTemplateDecl {{.*}} func2 // CHECK: FunctionDecl {{.*}} func2 'void ()' @@ -31,26 +23,6 @@ sycl::queue deviceQueue; template [[intel::scheduler_target_fmax_mhz(N)]] void func2() {} -template -[[intel::scheduler_target_fmax_mhz(N)]] void func3() {} - -// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored. -// CHECK: FunctionDecl {{.*}} {{.*}} func4 'void ()' -// CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} -// CHECK-NEXT: ConstantExpr {{.*}} 'int' -// CHECK-NEXT: value: Int 10 -// CHECK-NEXT: IntegerLiteral{{.*}}10{{$}} -[[intel::scheduler_target_fmax_mhz(10)]] -[[intel::scheduler_target_fmax_mhz(10)]] void -func4() {} - -class KernelFunctor { -public: - void operator()() const { - func1(); - } -}; - // Test that checks template parameter support on class member function. template class KernelFunctor2 { @@ -61,11 +33,6 @@ class KernelFunctor2 { int main() { deviceQueue.submit([&](sycl::handler &h) { - // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_1 - // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr - KernelFunctor f1; - h.single_task(f1); - // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_2 // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} // CHECK-NEXT: ConstantExpr{{.*}}'int' @@ -84,16 +51,6 @@ int main() { h.single_task( []() [[intel::scheduler_target_fmax_mhz(4)]]{}); - // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_4 - // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 75 - // CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} 'int' - // CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} referenced 'int' depth 0 index 0 N - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 75 - h.single_task( - []() { func3<75>(); }); - // Ignore duplicate attribute. h.single_task( // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_5 diff --git a/clang/test/SemaSYCL/sycl-2017-future-compat.cpp b/clang/test/SemaSYCL/sycl-2017-future-compat.cpp deleted file mode 100644 index 477c01367fc8..000000000000 --- a/clang/test/SemaSYCL/sycl-2017-future-compat.cpp +++ /dev/null @@ -1,13 +0,0 @@ -// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -sycl-std=2017 -verify %s -// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify=sycl-2017 %s - -// We do not expect any diagnostics from this file when disabling future compat -// warnings. -// sycl-2017-no-diagnostics - -// Test that we get compatibility warnings when using a SYCL 2020 attribute -// spelling while not in SYCL 2020 mode. -[[sycl::reqd_work_group_size(1, 1, 1)]] void f1(); // expected-warning {{use of attribute 'reqd_work_group_size' is a SYCL 2020 extension}} -[[sycl::work_group_size_hint(1, 1, 1)]] void f2(); // expected-warning {{use of attribute 'work_group_size_hint' is a SYCL 2020 extension}} -[[sycl::reqd_sub_group_size(1)]] void f3(); // expected-warning {{use of attribute 'reqd_sub_group_size' is a SYCL 2020 extension}} -[[sycl::vec_type_hint(int)]] void f4(); // expected-warning {{use of attribute 'vec_type_hint' is a SYCL 2020 extension}} diff --git a/clang/test/SemaSYCL/sycl-esimd-ast.cpp b/clang/test/SemaSYCL/sycl-esimd-ast.cpp index 54914f087c91..8c3e01aa529a 100644 --- a/clang/test/SemaSYCL/sycl-esimd-ast.cpp +++ b/clang/test/SemaSYCL/sycl-esimd-ast.cpp @@ -1,6 +1,6 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump %s | FileCheck %s -// Tests for AST of sycl_explicit_simd function attribute in SYCL 1.2.1. +// Tests for AST of sycl_explicit_simd function attribute in SYCL. #include "sycl.hpp" @@ -10,8 +10,6 @@ struct FuncObj { [[intel::sycl_explicit_simd]] void operator()() const {} }; -[[intel::sycl_explicit_simd]] void func() {} - int main() { deviceQueue.submit([&](sycl::handler &h) { // CHECK-LABEL: FunctionDecl {{.*}}test_kernel1 @@ -30,15 +28,6 @@ int main() { h.single_task( []() [[intel::sycl_explicit_simd]]{}); - // Test attribute is propagated. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 - // CHECK: SYCLSimdAttr {{.*}} Implicit - // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit - // CHECK-NEXT: AsmLabelAttr {{.*}} Implicit - // CHECK-NEXT: SYCLSimdAttr {{.*}} - // CHECK-NEXT: SYCLSimdAttr {{.*}} - h.single_task( - []() [[intel::sycl_explicit_simd]] { func(); }); }); return 0; } diff --git a/clang/test/SemaSYCL/sycl-esimd.cpp b/clang/test/SemaSYCL/sycl-esimd.cpp index 5153aa1ea3b4..349229c62559 100644 --- a/clang/test/SemaSYCL/sycl-esimd.cpp +++ b/clang/test/SemaSYCL/sycl-esimd.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify %s +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s // This test checks specifics of semantic analysis of ESIMD kernels. @@ -19,14 +19,6 @@ void kernel0(const F &f) __attribute__((sycl_kernel)) { f(); } -// expected-note@+1{{conflicting attribute is here}} -[[intel::reqd_sub_group_size(2)]] void g0() {} - -void test0() { - // expected-error@+2{{conflicting attributes applied to a SYCL kernel}} - // expected-note@+1{{conflicting attribute is here}} - kernel0([=]() __attribute__((sycl_explicit_simd)) { g0(); }); -} // -- Usual kernel can't call ESIMD function template @@ -34,13 +26,6 @@ void kernel1(const F &f) __attribute__((sycl_kernel)) { f(); } -// expected-note@+1{{attribute is here}} -__attribute__((sycl_explicit_simd)) void g1() {} - -void test1() { - // expected-error@+1{{SYCL kernel without 'sycl_explicit_simd' attribute can't call a function with this attribute}} - kernel1([=]() { g1(); }); -} // ----------- Positive tests diff --git a/clang/test/SemaSYCL/vec-type-hint.cpp b/clang/test/SemaSYCL/vec-type-hint.cpp index e9fccb4ae292..4a0dc4fe55b4 100644 --- a/clang/test/SemaSYCL/vec-type-hint.cpp +++ b/clang/test/SemaSYCL/vec-type-hint.cpp @@ -1,10 +1,8 @@ -// RUN: %clang_cc1 -fsycl-is-device -sycl-std=2017 -fsyntax-only -DKERNEL= -verify=sycl-2017 %s -// RUN: %clang_cc1 -fsycl-is-device -sycl-std=2020 -fsyntax-only -DKERNEL= -verify=sycl-2020 %s +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -DKERNEL= %s // RUN: %clang_cc1 -fsyntax-only -DKERNEL=kernel -verify=opencl -x cl %s // opencl-no-diagnostics -// sycl-2017-no-diagnostics -// __attribute__((vec_type_hint)) is deprecated without replacement in SYCL -// 2020 mode, but is allowed in SYCL 2017 and OpenCL modes. +// __attribute__((vec_type_hint)) is deprecated without replacement in SYCL 2020 mode, but +// is allowed in OpenCL mode. KERNEL __attribute__((vec_type_hint(int))) void foo() {} // sycl-2020-warning {{attribute 'vec_type_hint' is deprecated; attribute ignored}}