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 17a9d764b08dc..9a391a6cbba97 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 c6f7baf270541..d3f8fd40e3e6f 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 -sycl-std=2020 -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 97ab187974653..73107f27a5a36 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 -sycl-std=2020 -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 9845da905e0e2..d40a3d0e9110e 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 -sycl-std=2020 -triple spir64-unknown-unknown -o - | FileCheck %s template __attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { diff --git a/clang/test/CodeGenSYCL/launch_bounds_nvptx.cpp b/clang/test/CodeGenSYCL/launch_bounds_nvptx.cpp index 812363e227dd4..5bb3f34ea9947 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 -sycl-std=2020 -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, 8, 8), intel::min_work_groups_per_cu(N), - intel::max_work_groups_per_mp(N)]] void -zoo() {} - -[[intel::max_work_group_size(8, 8, 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 512} // CHECK: {{.*}}@{{.*}}kernel_name1, !"minnctapersm", i32 2} @@ -83,23 +66,9 @@ int main() { // CHECK: {{.*}}@{{.*}}Functor{{.*}}, !"maxntidx", i32 384} // CHECK: {{.*}}@{{.*}}Functor{{.*}}, !"minnctapersm", i32 6} // CHECK: {{.*}}@{{.*}}Functor{{.*}}, !"maxclusterrank", i32 6} -// CHECK: {{.*}}@{{.*}}kernel_name4, !"maxntidx", i32 512} -// CHECK: {{.*}}@{{.*}}kernel_name4, !"minnctapersm", i32 2} -// CHECK: {{.*}}@{{.*}}kernel_name4, !"maxclusterrank", i32 4} -// CHECK: {{.*}}@{{.*}}bar{{.*}}, !"maxntidx", i32 512} -// CHECK: {{.*}}@{{.*}}bar{{.*}}, !"minnctapersm", i32 2} -// CHECK: {{.*}}@{{.*}}bar{{.*}}, !"maxclusterrank", i32 4} -// CHECK: {{.*}}@{{.*}}kernel_name5, !"maxntidx", i32 1024} -// CHECK: {{.*}}@{{.*}}kernel_name5, !"minnctapersm", i32 16} -// CHECK: {{.*}}@{{.*}}kernel_name5, !"maxclusterrank", i32 16} -// CHECK: {{.*}}@{{.*}}zoo{{.*}}, !"maxntidx", i32 1024} -// CHECK: {{.*}}@{{.*}}zoo{{.*}}, !"minnctapersm", i32 16} -// CHECK: {{.*}}@{{.*}}zoo{{.*}}, !"maxclusterrank", i32 16} // CHECK: ![[MWGPC]] = !{i32 2} // CHECK: ![[MWGPM]] = !{i32 4} // CHECK: ![[MWGS]] = !{i32 8, i32 8, i32 8} // CHECK: ![[MWGPC_MWGPM]] = !{i32 6} // CHECK: ![[MWGS_2]] = !{i32 8, i32 8, i32 6} -// CHECK: ![[MWGPC_MWGPM_2]] = !{i32 16} -// CHECK: ![[MWGS_3]] = !{i32 8, i32 8, i32 16} diff --git a/clang/test/CodeGenSYCL/num-simd-work-items.cpp b/clang/test/CodeGenSYCL/num-simd-work-items.cpp index aeba63e0ffeea..ae2ecc651cac1 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 -sycl-std=2020 -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-work-group-size.cpp b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp index 2eff74dc8776f..4fb492543516a 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 -sycl-std=2020 -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=2020 -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=2020 -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=2020 -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 5194b92f44586..b8484e80944e2 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 -sycl-std=2020 -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 0e727251a1ce9..c8c523d4deceb 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 -sycl-std=2020 -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/SemaSYCL/intel-max-global-work-dim-device-ast.cpp b/clang/test/SemaSYCL/intel-max-global-work-dim-device-ast.cpp index 7ec419972dea2..5a2916faa53c5 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 -sycl-std=2020 -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/lb_sm_90_ast.cpp b/clang/test/SemaSYCL/lb_sm_90_ast.cpp index 2493445d6b985..3e6ea69d7a5a2 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 -sycl-std=2020 -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/scheduler_target_fmax_mhz_ast.cpp b/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp index eb6293280c041..c06ebb819db5e 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 -sycl-std=2020 -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