8000 Merge branch 'SYCLomatic' into texture · oneapi-src/SYCLomatic-test@c3e4c40 · GitHub
[go: up one dir, main page]
Skip to content

Commit c3e4c40

Browse files
ziranzhaziranzha
authored
Merge branch 'SYCLomatic' into texture
2 parents ec22465 + dae41b1 commit c3e4c40

File tree

5 files changed

+41
-55
lines changed

5 files changed

+41
-55
lines changed

‎features/config/TEMPLATE_math_intel_specific.xml

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4,6 +4,7 @@
44
<description>test</description>
55
<files>
66
<file path="feature_case/math/${testName}.cu" />
7+
<file path="feature_case/math/intel_specific_math.yaml" />
78
</files>
89
<rules>
910
<optlevelRule GPUFeature="NOT double" excludeOptlevelNameString="gpu" />

‎features/feature_case/math/intel_specific_math.yaml

Lines changed: 36 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,36 @@
1+
# Copyright (C) Intel Corporation
2+
# SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
3+
# See https://llvm.org/LICENSE.txt for license information.
4+
5+
# This file will be installed to folder:
6+
# {dpct_install_folder}/extensions/opt_rules.
7+
# The rule specified in this file can be imported with commandline option:
8+
# --rule-file={dpct_install_folder}/extensions/opt_rules/intel_specific_math.yaml
9+
---
10+
- Rule: intel_specific_math_erfinv
11+
Kind: API
12+
Priority: Takeover
13+
In: erfinv
14+
Out: sycl::ext::intel::math::erfinv($1)
15+
Includes: [<sycl/ext/intel/math.hpp>]
16+
17+
- Rule: intel_specific_math_erfinvf
18+
Kind: API
19+
Priority: Takeover
20+
In: erfinvf
21+
Out: sycl::ext::intel::math::erfinv($1)
22+
Includes: [<sycl/ext/intel/math.hpp>]
23+
24+
- Rule: intel_specific_math_normcdff
25+
Kind: API
26+
Priority: Takeover
27+
In: normcdff
28+
Out: sycl::ext::intel::math::cdfnorm($1)
29+
Includes: [<sycl/ext/intel/math.hpp>]
30+
31+
- Rule: intel_specific_math_normcdf
32+
Kind: API
33+
Priority: Takeover
34+
In: normcdf
35+
Out: sycl::ext::intel::math::cdfnorm($1)
36+
Includes: [<sycl/ext/intel/math.hpp>]

‎features/feature_case/math/math_intel_specific.cu

Lines changed: 1 addition & 53 deletions
Original file line numberDiff line numberDiff line change
@@ -7,58 +7,24 @@
77
//
88
// ===----------------------------------------------------------------------===//
99
// RUN: dpct --rule-file=%S/../../tools/dpct/DpctOptRules/intel_specific_math.yaml --format-range=none -out-root %T/math_specific_UDR_test %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only
10-
// RUN: FileCheck --input-file %T/math_specific_UDR_test/math_specific_UDR_test.dp.cpp --match-full-lines %s
11-
12-
// CHECK: #include <CL/sycl.hpp>
13-
// CHECK: #include <dpct/dpct.hpp>
14-
15-
// CHECK: #include <sycl/ext/intel/math.hpp>
1610

1711
#include "cuda_fp16.h"
1812
#include <iostream>
19-
// CHECK: void kernelFunc(double *deviceArray) {
20-
// CHECK: double &d0 = *deviceArray;
21-
// CHECK: d0 = sycl::ext::intel::math::erfinv(d0);
22-
// CHECK: d0 = sycl::ext::intel::math::cdfnorm(d0);
23-
// CHECK: }
13+
2414
__global__ void kernelFunc(double *deviceArray) {
2515
double &d0 = *deviceArray;
2616
d0 = erfinv(d0);
2717
d0 = normcdf(d0);
2818

2919
}
3020

31-
// CHECK: void kernelFunc(float *deviceArray) {
32-
// CHECK: float &f0 = *deviceArray;
33-
// CHECK: f0 = sycl::ext::intel::math::erfinv(f0);
34-
// CHECK: f0 = sycl::ext::intel::math::cdfnorm(f0);
35-
// CHECK: }
3621
__global__ void kernelFunc(float *deviceArray) {
3722
float &f0 = *deviceArray;
3823
f0 = erfinvf(f0);
3924
f0 = normcdff(f0);
4025

4126
}
4227

43-
44-
// CHECK: void testDouble() {
45-
// CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device();
46-
// CHECK: sycl::queue &q_ct1 = dev_ct1.default_queue();
47-
// CHECK: const unsigned int NUM = 1;
48-
// CHECK: const unsigned int bytes = NUM * sizeof(double);
49-
// CHECK: double *hostArrayDouble = (double *)malloc(bytes);
50-
// CHECK: memset(hostArrayDouble, 0, bytes);
51-
// CHECK: double *deviceArrayDouble;
52-
// CHECK: deviceArrayDouble = (double *)sycl::malloc_device(bytes, q_ct1);
53-
// CHECK: q_ct1.memcpy(deviceArrayDouble, hostArrayDouble, bytes).wait();
54-
// CHECK: q_ct1.parallel_for(
55-
// CHECK: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
56-
// CHECK: [=](sycl::nd_item<3> item_ct1) {
57-
// CHECK: kernelFunc(deviceArrayDouble);
58-
// CHECK: });
59-
// CHECK: q_ct1.memcpy(hostArrayDouble, deviceArrayDouble, bytes).wait();
60-
// CHECK: sycl::free(deviceArrayDouble, q_ct1);
61-
// CHECK: }
6228
bool testDouble() {
6329
double *hostArrayDouble = (double *)malloc(sizeof(double));
6430
*hostArrayDouble = 0.956841;
@@ -76,24 +42,6 @@ bool testDouble() {
7642
return true;
7743
}
7844

79-
// CHECK: void testFloat() {
80-
// CHECK: dpct::device_ext &dev_ct1 = dpct::get_current_device();
81-
// CHECK: sycl::queue &q_ct1 = dev_ct1.default_queue();
82-
// CHECK: const unsigned int NUM = 1;
83-
// CHECK: const unsigned int bytes = NUM * sizeof(float);
84-
// CHECK: float *hostArrayFloat = (float *)malloc(bytes);
85-
// CHECK: memset(hostArrayFloat, 0, bytes);
86-
// CHECK: float *deviceArrayFloat;
87-
// CHECK: deviceArrayFloat = (float *)sycl::malloc_device(bytes, q_ct1);
88-
// CHECK: q_ct1.memcpy(deviceArrayFloat, hostArrayFloat, bytes).wait();
89-
// CHECK: q_ct1.parallel_for(
90-
// CHECK: sycl::nd_range<3>(sycl::range<3>(1, 1, 1), sycl::range<3>(1, 1, 1)),
91-
// CHECK: [=](sycl::nd_item<3> item_ct1) {
92-
// CHECK: kernelFunc(deviceArrayFloat);
93-
// CHECK: });
94-
// CHECK: q_ct1.memcpy(hostArrayFloat, deviceArrayFloat, bytes).wait();
95-
// CHECK: sycl::free(deviceArrayFloat, q_ct1);
96-
// CHECK: }
9745
bool testFloat() {
9846
float *hostArrayFloat = (float *)malloc(sizeof(float));
9947
*hostArrayFloat = 0.1568541541f;

‎features/features.xml

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -359,5 +359,6 @@
359359
<test testName="libcu_atomic" configFile="config/TEMPLATE_libcu.xml" />
360360
<test testName="pointer_attributes" configFile="config/TEMPLATE_pointer_attributes.xml" />
361361
<test testName="image" configFile="config/TEMPLATE_image.xml" />
362+
<test testName="math_intel_specific" configFile="config/TEMPLATE_math_intel_specific.xml" />
362363
</tests>
363364
</suite>

‎features/test_feature.py

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -64,7 +64,7 @@ def migrate_test():
6464
if test_config.current_test in logical_group_exper:
6565
src.append(' --use-experimental-features=logical-group ')
6666
if test_config.current_test == 'math_intel_specific':
67-
src.append(' --rule-file=$(dirname $(which dpct))/../extensions/opt_rules/intel_specific_math.yaml ')
67+
src.append(' --rule-file=./math_intel_specific/intel_specific_math.yaml')
6868

6969
return do_migrate(src, in_root, test_config.out_root, extra_args)
7070

@@ -125,6 +125,6 @@ def run_test():
125125
return True
126126
os.environ['SYCL_DEVICE_FILTER'] = test_config.device_filter
127127
if test_config.current_test == 'ccl':
128-
return call_subprocess('mpirun -n 2 ' + os.path.join(os.path.curdir, test_config.current_test + '.run '));
128+
return call_subprocess('mpirun -n 2 ' + os.path.join(os.path.curdir, test_config.current_test + '.run '))
129129
return run_binary_with_args()
130130

0 commit comments

Comments
 (0)
0