Skip to content

[SYCLomatic] Fix the index base of the result of legacy cublasI<t>amin and cublasI<t>amax #2815

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
ziranzha merged 1 commit into from
May 6, 2025
Merged

[SYCLomatic] Fix the index base of the result of legacy cublasI<t>amin and cublasI<t>amax #2815

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
24 changes: 19 additions & 5 deletions clang/lib/DPCT/RulesMathLib/BLASAPIMigration.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -723,13 +723,20 @@ void BLASFunctionCallRule::runRule(const MatchFinder::MatchResult &Result) {
ReplInfo.BufferTypeInfo[ReplInfo.BufferTypeInfo.size() - 1];
std::string ReturnValueParamsStr;
if (DpctGlobalInfo::getUsmLevel() == UsmLevel::UL_Restricted) {
CallExprReplStr = CallExprReplStr + ", " + ResultTempPtr;
if (FuncName == "cublasIsamax" || FuncName == "cublasIdamax" ||
FuncName == "cublasIcamax" || FuncName == "cublasIzamax" ||
FuncName == "cublasIsamin" || FuncName == "cublasIdamin" ||
FuncName == "cublasIcamin" || FuncName == "cublasIzamin") {
CallExprReplStr = CallExprReplStr + ", oneapi::mkl::index_base::one";
}
requestFeature(HelperFeatureEnum::device_ext);
auto DefaultQueue = DpctGlobalInfo::getDefaultQueue(CE);
PrefixInsertStr = PrefixInsertStr + ResultType + "* " + ResultTempPtr +
" = " + MapNames::getClNamespace() +
"malloc_shared<" + ResultType + ">(1, " + DefaultQueue + ");" +
getNL() + IndentStr + CallExprReplStr + ", " +
ResultTempPtr + ").wait();" + getNL() + IndentStr;
"malloc_shared<" + ResultType + ">(1, " +
DefaultQueue + ");" + getNL() + IndentStr +
CallExprReplStr + ").wait();" + getNL() + IndentStr;

ReturnValueParamsStr =
"(" + ResultTempPtr + "->real(), " + ResultTempPtr + "->imag())";
Expand All @@ -748,11 +755,18 @@ void BLASFunctionCallRule::runRule(const MatchFinder::MatchResult &Result) {
ResultTempPtr + ", " + DefaultQueue + ");";
}
} else {
CallExprReplStr = CallExprReplStr + ", " + ResultTempBuf;
if (FuncName == "cublasIsamax" || FuncName == "cublasIdamax" ||
FuncName == "cublasIcamax" || FuncName == "cublasIzamax" ||
FuncName == "cublasIsamin" || FuncName == "cublasIdamin" ||
FuncName == "cublasIcamin" || FuncName == "cublasIzamin") {
CallExprReplStr = CallExprReplStr + ", oneapi::mkl::index_base::one";
}
PrefixInsertStr = PrefixInsertStr + MapNames::getClNamespace() +
"buffer<" + ResultType + "> " + ResultTempBuf + "(" +
MapNames::getClNamespace() + "range<1>(1));" +
getNL() + IndentStr + CallExprReplStr + ", " +
ResultTempBuf + ");" + getNL() + IndentStr;
getNL() + IndentStr + CallExprReplStr + ");" +
getNL() + IndentStr;
ReturnValueParamsStr =
"(" + ResultTempBuf + ".get_host_access(" +
MapNames::getClNamespace() + "read_only)[0].real(), " +
Expand Down
14 changes: 7 additions & 7 deletions clang/test/dpct/cublas-usm-legacy.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -66,30 +66,30 @@ int main() {

// CHECK: int res;
// CHECK-NEXT: int64_t* res_temp_ptr_ct{{[0-9]+}} = sycl::malloc_shared<int64_t>(1, q_ct1);
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, x_S, incx, res_temp_ptr_ct{{[0-9]+}}).wait();
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, x_S, incx, res_temp_ptr_ct{{[0-9]+}}, oneapi::mkl::index_base::one).wait();
// CHECK-NEXT: res = *res_temp_ptr_ct{{[0-9]+}};
// CHECK-NEXT:sycl::free(res_temp_ptr_ct{{[0-9]+}}, q_ct1);
int res = cublasIsamax(n, x_S, incx);
// CHECK: int64_t* res_temp_ptr_ct{{[0-9]+}} = sycl::malloc_shared<int64_t>(1, q_ct1);
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, x_D, incx, res_temp_ptr_ct{{[0-9]+}}).wait();
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, x_D, incx, res_temp_ptr_ct{{[0-9]+}}, oneapi::mkl::index_base::one).wait();
// CHECK-NEXT: res = *res_temp_ptr_ct{{[0-9]+}};
// CHECK-NEXT:sycl::free(res_temp_ptr_ct{{[0-9]+}}, q_ct1);
res = cublasIdamax(n, x_D, incx);
// CHECK: int64_t* res_temp_ptr_ct{{[0-9]+}} = sycl::malloc_shared<int64_t>(1, q_ct1);
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, (std::complex<float>*)x_C, incx, res_temp_ptr_ct{{[0-9]+}}).wait();
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, (std::complex<float>*)x_C, incx, res_temp_ptr_ct{{[0-9]+}}, oneapi::mkl::index_base::one).wait();
// CHECK-NEXT: res = *res_temp_ptr_ct{{[0-9]+}};
// CHECK-NEXT:sycl::free(res_temp_ptr_ct{{[0-9]+}}, q_ct1);
res = cublasIcamax(n, x_C, incx);
// CHECK: int64_t* res_temp_ptr_ct{{[0-9]+}} = sycl::malloc_shared<int64_t>(1, q_ct1);
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, (std::complex<double>*)x_Z, incx, res_temp_ptr_ct{{[0-9]+}}).wait();
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, (std::complex<double>*)x_Z, incx, res_temp_ptr_ct{{[0-9]+}}, oneapi::mkl::index_base::one).wait();
// CHECK-NEXT: res = *res_temp_ptr_ct{{[0-9]+}};
// CHECK-NEXT:sycl::free(res_temp_ptr_ct{{[0-9]+}}, q_ct1);
res = cublasIzamax(n, x_Z, incx);

// Because the return value of origin API is the result value, not the status, so keep using lambda here.
// CHECK: if([&](){
// CHECK-NEXT: int64_t* res_temp_ptr_ct{{[0-9]+}} = sycl::malloc_shared<int64_t>(1, q_ct1);
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, (std::complex<double>*)x_Z, incx, res_temp_ptr_ct{{[0-9]+}}).wait();
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, (std::complex<double>*)x_Z, incx, res_temp_ptr_ct{{[0-9]+}}, oneapi::mkl::index_base::one).wait();
// CHECK-NEXT: int64_t res_temp_val_ct{{[0-9]+}} = *res_temp_ptr_ct{{[0-9]+}};
// CHECK-NEXT: sycl::free(res_temp_ptr_ct{{[0-9]+}}, q_ct1);
// CHECK-NEXT: return res_temp_val_ct{{[0-9]+}};
Expand All @@ -98,7 +98,7 @@ int main() {

// CHECK: if(0!=[&](){
// CHECK-NEXT: int64_t* res_temp_ptr_ct{{[0-9]+}} = sycl::malloc_shared<int64_t>(1, q_ct1);
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, (std::complex<double>*)x_Z, incx, res_temp_ptr_ct{{[0-9]+}}).wait();
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, (std::complex<double>*)x_Z, incx, res_temp_ptr_ct{{[0-9]+}}, oneapi::mkl::index_base::one).wait();
// CHECK-NEXT: int64_t res_temp_val_ct{{[0-9]+}} = *res_temp_ptr_ct{{[0-9]+}};
// CHECK-NEXT: sycl::free(res_temp_ptr_ct{{[0-9]+}}, q_ct1);
// CHECK-NEXT: return res_temp_val_ct{{[0-9]+}};
Expand Down Expand Up @@ -233,7 +233,7 @@ int main() {
//CHECK:int foo(){
//CHECK-NEXT: return [&](){
//CHECK-NEXT: int64_t* res_temp_ptr_ct{{[0-9]+}} = sycl::malloc_shared<int64_t>(1, dpct::get_in_order_queue());
//CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, (std::complex<double>*)x_Z, incx, res_temp_ptr_ct{{[0-9]+}}).wait();
//CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, (std::complex<double>*)x_Z, incx, res_temp_ptr_ct{{[0-9]+}}, oneapi::mkl::index_base::one).wait();
//CHECK-NEXT: int64_t res_temp_val_ct{{[0-9]+}} = *res_temp_ptr_ct{{[0-9]+}};
//CHECK-NEXT: sycl::free(res_temp_ptr_ct{{[0-9]+}}, dpct::get_in_order_queue());
//CHECK-NEXT: return res_temp_val_ct{{[0-9]+}};
Expand Down
8 changes: 4 additions & 4 deletions clang/test/dpct/cublasLegacyCZ.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -60,15 +60,15 @@ int main() {
// CHECK-NEXT: {
// CHECK-NEXT: auto x_C_buf_ct{{[0-9]+}} = dpct::get_buffer<std::complex<float>>(x_C);
// CHECK-NEXT: sycl::buffer<int64_t> res_temp_buf_ct{{[0-9]+}}(sycl::range<1>(1));
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, x_C_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}});
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, x_C_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}}, oneapi::mkl::index_base::one);
// CHECK-NEXT: res = res_temp_buf_ct{{[0-9]+}}.get_host_access(sycl::read_only)[0];
// CHECK-NEXT: }
int res = cublasIcamax(n, x_C, incx);

// CHECK: {
// CHECK-NEXT: auto x_Z_buf_ct{{[0-9]+}} = dpct::get_buffer<std::complex<double>>(x_Z);
// CHECK-NEXT: sycl::buffer<int64_t> res_temp_buf_ct{{[0-9]+}}(sycl::range<1>(1));
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, x_Z_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}});
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, x_Z_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}}, oneapi::mkl::index_base::one);
// CHECK-NEXT: *result = res_temp_buf_ct{{[0-9]+}}.get_host_access(sycl::read_only)[0];
// CHECK-NEXT: }
*result = cublasIzamax(n, x_Z, incx);
Expand All @@ -77,15 +77,15 @@ int main() {
// CHECK: {
// CHECK-NEXT: auto x_C_buf_ct{{[0-9]+}} = dpct::get_buffer<std::complex<float>>(x_C);
// CHECK-NEXT: sycl::buffer<int64_t> res_temp_buf_ct{{[0-9]+}}(sycl::range<1>(1));
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamin(dpct::blas::descriptor::get_saved_queue(), n, x_C_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}});
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamin(dpct::blas::descriptor::get_saved_queue(), n, x_C_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}}, oneapi::mkl::index_base::one);
// CHECK-NEXT: *result = res_temp_buf_ct{{[0-9]+}}.get_host_access(sycl::read_only)[0];
// CHECK-NEXT: }
*result = cublasIcamin(n, x_C, incx);

// CHECK: {
// CHECK-NEXT: auto x_Z_buf_ct{{[0-9]+}} = dpct::get_buffer<std::complex<double>>(x_Z);
// CHECK-NEXT: sycl::buffer<int64_t> res_temp_buf_ct{{[0-9]+}}(sycl::range<1>(1));
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamin(dpct::blas::descriptor::get_saved_queue(), n, x_Z_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}});
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamin(dpct::blas::descriptor::get_saved_queue(), n, x_Z_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}}, oneapi::mkl::index_base::one);
// CHECK-NEXT: *result = res_temp_buf_ct{{[0-9]+}}.get_host_access(sycl::read_only)[0];
// CHECK-NEXT: }
*result = cublasIzamin(n, x_Z, incx);
Expand Down
12 changes: 6 additions & 6 deletions clang/test/dpct/cublasLegacyLv123.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,15 +50,15 @@ int main() {
// CHECK-NEXT: {
// CHECK-NEXT: auto x_S_buf_ct{{[0-9]+}} = dpct::get_buffer<float>(x_S);
// CHECK-NEXT: sycl::buffer<int64_t> res_temp_buf_ct{{[0-9]+}}(sycl::range<1>(1));
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, x_S_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}});
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, x_S_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}}, oneapi::mkl::index_base::one);
// CHECK-NEXT: res = res_temp_buf_ct{{[0-9]+}}.get_host_access(sycl::read_only)[0];
// CHECK-NEXT: }
int res = cublasIsamax(n, x_S, incx);

// CHECK: {
// CHECK-NEXT: auto x_D_buf_ct{{[0-9]+}} = dpct::get_buffer<double>(x_D);
// CHECK-NEXT: sycl::buffer<int64_t> res_temp_buf_ct{{[0-9]+}}(sycl::range<1>(1));
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, x_D_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}});
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, x_D_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}}, oneapi::mkl::index_base::one);
// CHECK-NEXT: *result = res_temp_buf_ct{{[0-9]+}}.get_host_access(sycl::read_only)[0];
// CHECK-NEXT: }
*result = cublasIdamax(n, x_D, incx);
Expand All @@ -67,15 +67,15 @@ int main() {
// CHECK: {
// CHECK-NEXT: auto x_S_buf_ct{{[0-9]+}} = dpct::get_buffer<float>(x_S);
// CHECK-NEXT: sycl::buffer<int64_t> res_temp_buf_ct{{[0-9]+}}(sycl::range<1>(1));
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamin(dpct::blas::descriptor::get_saved_queue(), n, x_S_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}});
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamin(dpct::blas::descriptor::get_saved_queue(), n, x_S_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}}, oneapi::mkl::index_base::one);
// CHECK-NEXT: *result = res_temp_buf_ct{{[0-9]+}}.get_host_access(sycl::read_only)[0];
// CHECK-NEXT: }
*result = cublasIsamin(n, x_S, incx);

// CHECK: {
// CHECK-NEXT: auto x_D_buf_ct{{[0-9]+}} = dpct::get_buffer<double>(x_D);
// CHECK-NEXT: sycl::buffer<int64_t> res_temp_buf_ct{{[0-9]+}}(sycl::range<1>(1));
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamin(dpct::blas::descriptor::get_saved_queue(), n, x_D_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}});
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamin(dpct::blas::descriptor::get_saved_queue(), n, x_D_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}}, oneapi::mkl::index_base::one);
// CHECK-NEXT: *result = res_temp_buf_ct{{[0-9]+}}.get_host_access(sycl::read_only)[0];
// CHECK-NEXT: }
*result = cublasIdamin(n, x_D, incx);
Expand Down Expand Up @@ -627,7 +627,7 @@ int main() {
// CHECK: for(int i = [&](){
// CHECK-NEXT: auto x_S_buf_ct{{[0-9]+}} = dpct::get_buffer<float>(x_S);
// CHECK-NEXT: sycl::buffer<int64_t> res_temp_buf_ct{{[0-9]+}}(sycl::range<1>(1));
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, x_S_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}});
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, x_S_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}}, oneapi::mkl::index_base::one);
// CHECK-NEXT: return res_temp_buf_ct{{[0-9]+}}.get_host_access(sycl::read_only)[0];
// CHECK-NEXT: }();;){}
for(int i = cublasIsamax(n, x_S, incx);;){}
Expand All @@ -640,7 +640,7 @@ int main() {
//CHECK-NEXT: return [&](){
//CHECK-NEXT: auto x_S_buf_ct{{[0-9]+}} = dpct::get_buffer<float>(x_S);
//CHECK-NEXT: sycl::buffer<int64_t> res_temp_buf_ct{{[0-9]+}}(sycl::range<1>(1));
//CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, x_S_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}});
//CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(dpct::blas::descriptor::get_saved_queue(), n, x_S_buf_ct{{[0-9]+}}, incx, res_temp_buf_ct{{[0-9]+}}, oneapi::mkl::index_base::one);
//CHECK-NEXT: return res_temp_buf_ct{{[0-9]+}}.get_host_access(sycl::read_only)[0];
//CHECK-NEXT: }();
//CHECK-NEXT:}
Expand Down
2 changes: 1 addition & 1 deletion clang/test/dpct/error-handling.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -705,7 +705,7 @@ void foo12() {
// CHECK-NEXT: sycl::buffer<int64_t> res_temp_buf_ct{{[0-9]+}}(sycl::range<1>(1));
// CHECK-NEXT: oneapi::mkl::blas::column_major::iamax(
// CHECK-NEXT: dpct::blas::descriptor::get_saved_queue(), 10, ct_0_buf_ct{{[0-9]+}}, 0,
// CHECK-NEXT: res_temp_buf_ct{{[0-9]+}});
// CHECK-NEXT: res_temp_buf_ct{{[0-9]+}}, oneapi::mkl::index_base::one);
// CHECK-NEXT: res = res_temp_buf_ct{{[0-9]+}}.get_host_access(sycl::read_only)[0];
// CHECK-NEXT: }
// CHECK-NEXT: }
Expand Down