Skip to content

Commit 8f7b5d2

Browse files
rbegamfadeeval
rbegam
authored and
fadeeval
committed
[SYCL] Fix incorrect data type (#2394)
Signed-off-by: rbegam <[email protected]>
1 parent eb14539 commit 8f7b5d2

12 files changed

+396
-0
lines changed

‎sycl/test/inline-asm/negative_tests/asm_bad_opcode.cpp

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,33 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -DINLINE_ASM -o %t.out
4+
// RUN: not %t.out 2>&1 | FileCheck %s
5+
6+
// CHECK: syntax error
7+
8+
#include "../include/asmhelper.h"
9+
#include <CL/sycl.hpp>
10+
11+
using dataType = cl::sycl::cl_int;
12+
13+
template <typename T = dataType>
14+
struct KernelFunctor : WithOutputBuffer<T> {
15+
KernelFunctor(size_t problem_size) : WithOutputBuffer<T>(problem_size) {}
16+
17+
void operator()(cl::sycl::handler &cgh) {
18+
cgh.parallel_for<KernelFunctor<T>>(
19+
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
20+
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
21+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
22+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
23+
"movi (M1_NM, 8) tmp1(0,1)<1> tmp2(0,0)\n");
24+
#endif
25+
});
26+
}
27+
};
28+
29+
int main() {
30+
KernelFunctor<> f(DEFAULT_PROBLEM_SIZE);
31+
launchInlineASMTest(f);
32+
return 0;
33+
}

‎sycl/test/inline-asm/negative_tests/asm_bad_operand_syntax.cpp

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,33 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -DINLINE_ASM -o %t.out
4+
// RUN: not %t.out 2>&1 | FileCheck %s
5+
6+
// CHECK: syntax error
7+
8+
#include "../include/asmhelper.h"
9+
#include <CL/sycl.hpp>
10+
11+
using dataType = cl::sycl::cl_int;
12+
13+
template <typename T = dataType>
14+
struct KernelFunctor : WithOutputBuffer<T> {
15+
KernelFunctor(size_t problem_size) : WithOutputBuffer<T>(problem_size) {}
16+
17+
void operator()(cl::sycl::handler &cgh) {
18+
cgh.parallel_for<KernelFunctor<T>>(
19+
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
20+
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
21+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
22+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
23+
"mov (M1_NM, 8) tmp1(0,1)<1>:f tmp2(0,0)<1;1,0>\n");
24+
#endif
25+
});
26+
}
27+
};
28+
29+
int main() {
30+
KernelFunctor<> f(DEFAULT_PROBLEM_SIZE);
31+
launchInlineASMTest(f);
32+
return 0;
33+
}

‎sycl/test/inline-asm/negative_tests/asm_duplicate_label.cpp

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,33 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -DINLINE_ASM -o %t.out
4+
// RUN: not %t.out 2>&1 | FileCheck %s
5+
6+
// CHECK: label is redefined
7+
8+
#include "../include/asmhelper.h"
9+
#include <CL/sycl.hpp>
10+
11+
using dataType = cl::sycl::cl_int;
12+
13+
template <typename T = dataType>
14+
struct KernelFunctor : WithOutputBuffer<T> {
15+
KernelFunctor(size_t problem_size) : WithOutputBuffer<T>(problem_size) {}
16+
17+
void operator()(cl::sycl::handler &cgh) {
18+
cgh.parallel_for<KernelFunctor<T>>(
19+
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
20+
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
21+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
22+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
23+
"check_label0:\\ncheck_label0:\n");
24+
#endif
25+
});
26+
}
27+
};
28+
29+
int main() {
30+
KernelFunctor<> f(DEFAULT_PROBLEM_SIZE);
31+
launchInlineASMTest(f);
32+
return 0;
33+
}

‎sycl/test/inline-asm/negative_tests/asm_illegal_exec_size.cpp

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,33 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -DINLINE_ASM -o %t.out
4+
// RUN: not %t.out 2>&1 | FileCheck %s
5+
6+
// CHECK: invalid execution size
7+
8+
#include "../include/asmhelper.h"
9+
#include <CL/sycl.hpp>
10+
11+
using dataType = cl::sycl::cl_int;
12+
13+
template <typename T = dataType>
14+
struct KernelFunctor : WithOutputBuffer<T> {
15+
KernelFunctor(size_t problem_size) : WithOutputBuffer<T>(problem_size) {}
16+
17+
void operator()(cl::sycl::handler &cgh) {
18+
cgh.parallel_for<KernelFunctor<T>>(
19+
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
20+
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
21+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
22+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
23+
"mov (M1_NM, 6) tmp1(0,1)<1> tmp2(0,0)<1;1,0>\n");
24+
#endif
25+
});
26+
}
27+
};
28+
29+
int main() {
30+
KernelFunctor<> f(DEFAULT_PROBLEM_SIZE);
31+
launchInlineASMTest(f);
32+
return 0;
33+
}

‎sycl/test/inline-asm/negative_tests/asm_missing_label.cpp

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,33 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -DINLINE_ASM -o %t.out
4+
// RUN: not %t.out 2>&1 | FileCheck %s
5+
6+
// CHECK: undefined label: check_label0
7+
8+
#include "../include/asmhelper.h"
9+
#include <CL/sycl.hpp>
10+
11+
using dataType = cl::sycl::cl_int;
12+
13+
template <typename T = dataType>
14+
struct KernelFunctor : WithOutputBuffer<T> {
15+
KernelFunctor(size_t problem_size) : WithOutputBuffer<T>(problem_size) {}
16+
17+
void operator()(cl::sycl::handler &cgh) {
18+
cgh.parallel_for<KernelFunctor<T>>(
19+
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
20+
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
21+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
22+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
23+
"goto (M1, 16) check_label0\n");
24+
#endif
25+
});
26+
}
27+
};
28+
29+
int main() {
30+
KernelFunctor<> f(DEFAULT_PROBLEM_SIZE);
31+
launchInlineASMTest(f);
32+
return 0;
33+
}

‎sycl/test/inline-asm/negative_tests/asm_missing_region.cpp

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,33 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -DINLINE_ASM -o %t.out
4+
// RUN: not %t.out 2>&1 | FileCheck %s
5+
6+
// CHECK: syntax error
7+
8+
#include "../include/asmhelper.h"
9+
#include <CL/sycl.hpp>
10+
11+
using dataType = cl::sycl::cl_int;
12+
13+
template <typename T = dataType>
14+
struct KernelFunctor : WithOutputBuffer<T> {
15+
KernelFunctor(size_t problem_size) : WithOutputBuffer<T>(problem_size) {}
16+
17+
void operator()(cl::sycl::handler &cgh) {
18+
cgh.parallel_for<KernelFunctor<T>>(
19+
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
20+
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
21+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
22+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
23+
"mov (M1_NM, 8) tmp1(0,1)<1> tmp2(0,0)\n");
24+
#endif
25+
});
26+
}
27+
};
28+
29+
int main() {
30+
KernelFunctor<> f(DEFAULT_PROBLEM_SIZE);
31+
launchInlineASMTest(f);
32+
return 0;
33+
}

‎sycl/test/inline-asm/negative_tests/asm_simple.cpp

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,33 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -DINLINE_ASM -o %t.out
4+
// RUN: not %t.out 2>&1 | FileCheck %s
5+
6+
// CHECK: syntax error
7+
8+
#include "../include/asmhelper.h"
9+
#include <CL/sycl.hpp>
10+
11+
using dataType = cl::sycl::cl_int;
12+
13+
template <typename T = dataType>
14+
struct KernelFunctor : WithOutputBuffer<T> {
15+
KernelFunctor(size_t problem_size) : WithOutputBuffer<T>(problem_size) {}
16+
17+
void operator()(cl::sycl::handler &cgh) {
18+
cgh.parallel_for<KernelFunctor<T>>(
19+
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
20+
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
21+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
22+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
23+
"@@\n");
24+
#endif
25+
});
26+
}
27+
};
28+
29+
int main() {
30+
KernelFunctor<> f(DEFAULT_PROBLEM_SIZE);
31+
launchInlineASMTest(f);
32+
return 0;
33+
}

‎sycl/test/inline-asm/negative_tests/asm_undefined_decl.cpp

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,33 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -DINLINE_ASM -o %t.out
4+
// RUN: not %t.out 2>&1 | FileCheck %s
5+
6+
// CHECK: syntax error
7+
8+
#include "../include/asmhelper.h"
9+
#include <CL/sycl.hpp>
10+
11+
using dataType = cl::sycl::cl_int;
12+
13+
template <typename T = dataType>
14+
struct KernelFunctor : WithOutputBuffer<T> {
15+
KernelFunctor(size_t problem_size) : WithOutputBuffer<T>(problem_size) {}
16+
17+
void operator()(cl::sycl::handler &cgh) {
18+
cgh.parallel_for<KernelFunctor<T>>(
19+
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
20+
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
21+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
22+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
23+
"mov (M1_NM, 8) tmp1(0,1)<1> my_super_var(0,0)\n");
24+
#endif
25+
});
26+
}
27+
};
28+
29+
int main() {
30+
KernelFunctor<> f(DEFAULT_PROBLEM_SIZE);
31+
launchInlineASMTest(f);
32+
return 0;
33+
}

‎sycl/test/inline-asm/negative_tests/asm_undefined_pred.cpp

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,33 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -DINLINE_ASM -o %t.out
4+
// RUN: not %t.out 2>&1 | FileCheck %s
5+
6+
// CHECK: undefined predicate variable
7+
8+
#include "../include/asmhelper.h"
9+
#include <CL/sycl.hpp>
10+
11+
using dataType = cl::sycl::cl_int;
12+
13+
template <typename T = dataType>
14+
struct KernelFunctor : WithOutputBuffer<T> {
15+
KernelFunctor(size_t problem_size) : WithOutputBuffer<T>(problem_size) {}
16+
17+
void operator()(cl::sycl::handler &cgh) {
18+
cgh.parallel_for<KernelFunctor<T>>(
19+
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
20+
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
21+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
22+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
23+
"cmp.lt (M1_NM, 8) P3 tmp1(0,0)<0;1,0> 0x3:ud\n");
24+
#endif
25+
});
26+
}
27+
};
28+
29+
int main() {
30+
KernelFunctor<> f(DEFAULT_PROBLEM_SIZE);
31+
launchInlineASMTest(f);
32+
return 0;
33+
}

‎sycl/test/inline-asm/negative_tests/asm_unsup_instruction_divm_skl.cpp

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,33 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -DINLINE_ASM -o %t.out
4+
// RUN: not %t.out 2>&1 | FileCheck %s
5+
6+
// CHECK: Build succeeded
7+
8+
#include "../include/asmhelper.h"
9+
#include <CL/sycl.hpp>
10+
11+
using dataType = cl::sycl::cl_int;
12+
13+
template <typename T = dataType>
14+
struct KernelFunctor : WithOutputBuffer<T> {
15+
KernelFunctor(size_t problem_size) : WithOutputBuffer<T>(problem_size) {}
16+
17+
void operator()(cl::sycl::handler &cgh) {
18+
cgh.parallel_for<KernelFunctor<T>>(
19+
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
20+
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
21+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
22+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
23+
"divm (M1, 8) tmp1(0,0)<1> tmp2(0,0)<1;1,0> 0x2:f\n");
24+
#endif
25+
});
26+
}
27+
};
28+
29+
int main() {
30+
KernelFunctor<> f(DEFAULT_PROBLEM_SIZE);
31+
launchInlineASMTest(f);
32+
return 0;
33+
}

‎sycl/test/inline-asm/negative_tests/asm_unsup_instruction_rol_skl.cpp

Lines changed: 33 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,33 @@
1+
// UNSUPPORTED: cuda
2+
// REQUIRES: gpu,linux
3+
// RUN: %clangxx -fsycl %s -DINLINE_ASM -o %t.out
4+
// RUN: not %t.out 2>&1 | FileCheck %s
5+
6+
// CHECK: Build succeeded
7+
8+
#include "../include/asmhelper.h"
9+
#include <CL/sycl.hpp>
10+
11+
using dataType = cl::sycl::cl_int;
12+
13+
template <typename T = dataType>
14+
struct KernelFunctor : WithOutputBuffer<T> {
15+
KernelFunctor(size_t problem_size) : WithOutputBuffer<T>(problem_size) {}
16+
17+
void operator()(cl::sycl::handler &cgh) {
18+
cgh.parallel_for<KernelFunctor<T>>(
19+
cl::sycl::range<1>{this->getOutputBufferSize()}, [=](cl::sycl::id<1> wiID) [[cl::intel_reqd_sub_group_size(8)]] {
20+
#if defined(INLINE_ASM) && defined(__SYCL_DEVICE_ONLY__)
21+
asm volatile(".decl tmp1 v_type=G type=d num_elts=16 align=GRF\n"
22+
".decl tmp2 v_type=G type=d num_elts=16 align=GRF\n"
23+
"rol (M1, 8) tmp1(0,0)<1> tmp2(0,0)<1;1,0> 0x2:d\n");
24+
#endif
25+
});
26+
}
27+
};
28+
29+
int main() {
30+
KernelFunctor<> f(DEFAULT_PROBLEM_SIZE);
31+
launchInlineASMTest(f);
32+
return 0;
33+
}

0 commit comments

Comments
 (0)