Skip to content

Commit

Permalink
[SYCL] getting rid of test cases based on attribute propageion via de…
Browse files Browse the repository at this point in the history 
…vice-functions incompatible with -sycl-std=2020
  • Loading branch information
@AndreiZibrov
AndreiZibrov committed Jun 21, 2024
1 parent cdd55bc commit 85d1860
Show file tree
Hide file tree
Showing 12 changed files with 16 additions and 277 deletions.
10 changes: 1 addition & 9 deletions clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp
View file
Original file line number Diff line number Diff line change
@@ -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"

Expand All @@ -16,9 +16,6 @@ class Functor {
[[intel::no_global_work_offset(SIZE)]] void operator()() const {}
};

template <int N>
[[intel::no_global_work_offset(N)]] void func() {}

int main() {
q.submit([&](handler &h) {
Foo boo;
Expand All @@ -32,10 +29,6 @@ int main() {

Functor<1> f;
h.single_task<class kernel_name4>(f);

h.single_task<class kernel_name5>([]() {
func<1>();
});
});
return 0;
}
Expand All @@ -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]] = !{}
10 changes: 1 addition & 9 deletions clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp
View file
Original file line number Diff line number Diff line change
@@ -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"

Expand All @@ -16,9 +16,6 @@ class Functor {
[[intel::max_global_work_dim(SIZE)]] void operator()() const {}
};

template <int N>
[[intel::max_global_work_dim(N)]] void func() {}

int main() {
q.submit([&](handler &h) {
Foo boo;
Expand All @@ -29,17 +26,12 @@ int main() {

Functor<2> f;
h.single_task<class kernel_name3>(f);

h.single_task<class kernel_name4>([]() {
func<2>();
});
});
return 0;
}

// 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}
9 changes: 1 addition & 8 deletions clang/test/CodeGenSYCL/intel-max-work-group-size.cpp
View file
Original file line number Diff line number Diff line change
@@ -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"

Expand All @@ -21,8 +21,6 @@ class Functor {
[[intel::max_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() const {}
};

template <int N, int N1, int N2>
[[intel::max_work_group_size(N, N1, N2)]] void func() {}

int main() {
q.submit([&](handler &h) {
Expand All @@ -38,9 +36,6 @@ int main() {
Functor<2, 2, 2> f;
h.single_task<class kernel_name4>(f);

h.single_task<class kernel_name5>([]() {
func<4, 4, 4>();
});
});
return 0;
}
Expand All @@ -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}
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/intel-restrict.cpp
View file
Original file line number Diff line number Diff line change
@@ -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 <typename name, typename Func>
__attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) {
Expand Down
33 changes: 1 addition & 32 deletions clang/test/CodeGenSYCL/launch_bounds_nvptx.cpp
View file
Original file line number Diff line number Diff line change
@@ -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
Expand All @@ -25,15 +25,6 @@ template <int N> class Functor {
operator()() const {}
};

template <int N>
[[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.
Expand All @@ -49,21 +40,13 @@ int main() {
// Test class template argument.
Functor<6> f;
h.single_task<class kernel_name3>(f);

// Test attribute is propagated.
h.single_task<class kernel_name4>([]() { bar(); });

// Test function template argument.
h.single_task<class kernel_name5>([]() { zoo<16>(); });
});
return 0;
}

// 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}
Expand All @@ -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}
9 changes: 1 addition & 8 deletions clang/test/CodeGenSYCL/num-simd-work-items.cpp
View file
Original file line number Diff line number Diff line change
@@ -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"

Expand All @@ -16,8 +16,6 @@ class Functor {
[[intel::num_simd_work_items(SIZE)]] void operator()() const {}
};

template <int N>
[[intel::num_simd_work_items(N)]] void func() {}

int main() {
q.submit([&](handler &h) {
Expand All @@ -30,18 +28,13 @@ int main() {
Functor<2> f;
h.single_task<class kernel_name3>(f);

h.single_task<class kernel_name4>([]() {
func<4>();
});
});
return 0;
}

// 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}
82 changes: 5 additions & 77 deletions clang/test/CodeGenSYCL/reqd-work-group-size.cpp
View file
Original file line number Diff line number Diff line change
@@ -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"

Expand All @@ -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 <int SIZE>
class FunctorTemp1D {
Expand All @@ -76,134 +47,91 @@ class CLFunctorTemp3D {
[[cl::reqd_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() const {}
};

template <int N>
[[sycl::reqd_work_group_size(N)]] void func1D() {}
template <int N, int N1>
[[sycl::reqd_work_group_size(N, N1)]] void func2D() {}
template <int N, int N1, int N2>
[[sycl::reqd_work_group_size(N, N1, N2)]] void func3D() {}
template <int N, int N1, int N2>
[[cl::reqd_work_group_size(N, N1, N2)]] void clfunc3D() {}

int main() {
q.submit([&](handler &h) {
CLFunctor32x16x16 clf32x16x16;
h.single_task<class kernel_name1>(clf32x16x16);

CLFunctor3D clf3d;
h.single_task<class kernel_name2>(clf3d);


h.single_task<class kernel_name3>(
[]() [[cl::reqd_work_group_size(8, 8, 8)]]{});

CLFunctorTemp3D<2, 2, 2> clft3d;
h.single_task<class kernel_name4>(clft3d);

h.single_task<class kernel_name5>([]() {
clfunc3D<8, 4, 4>();
});

h.single_task<class kernel_name6>(
[]() [[cl::reqd_work_group_size(1, 8, 2)]]{});

Functor32x16x16 f32x16x16;
h.single_task<class kernel_name7>(f32x16x16);

Functor3D f3d;
h.single_task<class kernel_name8>(f3d);

h.single_task<class kernel_name9>(
[]() [[sycl::reqd_work_group_size(8, 8, 8)]]{});

FunctorTemp3D<2, 2, 2> ft3d;
h.single_task<class kernel_name10>(ft3d);

h.single_task<class kernel_name11>([]() {
func3D<8, 4, 4>();
});

h.single_task<class kernel_name12>(
[]() [[sycl::reqd_work_group_size(1, 8, 2)]]{});

Functor32x16 f32x16;
h.single_task<class kernel_name13>(f32x16);

Functor2D f2d;
h.single_task<class kernel_name14>(f2d);

h.single_task<class kernel_name15>(
[]() [[sycl::reqd_work_group_size(8, 8)]]{});

FunctorTemp2D<2, 2> ft2d;
h.single_task<class kernel_name16>(ft2d);

h.single_task<class kernel_name17>([]() {
func2D<8, 4>();
});

h.single_task<class kernel_name18>(
[]() [[sycl::reqd_work_group_size(1, 8)]]{});

Functor32 f32;
h.single_task<class kernel_name19>(f32);

Functor1D f1d;
h.single_task<class kernel_name20>(f1d);

h.single_task<class kernel_name21>(
[]() [[sycl::reqd_work_group_size(8)]]{});

FunctorTemp1D<2> ft1d;
h.single_task<class kernel_name22>(ft1d);

h.single_task<class kernel_name23>([]() {
func1D<8>();
});

h.single_task<class kernel_name24>(
[]() [[sycl::reqd_work_group_size(1)]]{});
});
return 0;
}

// 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}
Expand Down
Loading

0 comments on commit 85d1860

Please sign in to comment.