Skip to content

[SYCLomatic] Fix migration issue for free-function-queries #2804

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 6 commits into from
Apr 30, 2025
Merged

[SYCLomatic] Fix migration issue for free-function-queries #2804

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
1ece36c
[SYCLomatic] Fix migration issue for free-function-queries
zhiweij1 Apr 24, 2025
aec62ae
Refine
zhiweij1 Apr 28, 2025
3f415dc
Merge remote-tracking branch 'origin/SYCLomatic' into free_func
zhiweij1 Apr 28, 2025
22acd3b
Fix
zhiweij1 Apr 28, 2025
085ca92
Fix
zhiweij1 Apr 29, 2025
df11eb8
Merge remote-tracking branch 'origin/SYCLomatic' into free_func
zhiweij1 Apr 29, 2025
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
44 changes: 27 additions & 17 deletions clang/lib/DPCT/AnalysisInfo.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -326,16 +326,22 @@ class FreeQueriesInfo {

static const FreeQueriesNames &getNames(FreeQueriesKind);
static std::shared_ptr<FreeQueriesInfo> getInfo(const FunctionDecl *);
static void printFreeQueriesFunctionName(llvm::raw_ostream &OS,
FreeQueriesKind K,
unsigned Dimension) {
template <typename T>
static typename std::enable_if<std::is_same_v<T, unsigned> ||
std::is_same_v<T, std::string>>::type
printFreeQueriesFunctionName(llvm::raw_ostream &OS, FreeQueriesKind K,
T Dimension) {
OS << getNames(K).FreeQueriesFuncName;
if (K != FreeQueriesKind::SubGroup) {
OS << '<';
if (Dimension) {
OS << Dimension;
if constexpr (std::is_same_v<T, unsigned>) {
if (Dimension) {
OS << Dimension;
} else {
OS << "dpct_placeholder /* Fix the dimension manually */";
}
} else {
OS << "dpct_placeholder /* Fix the dimension manually */";
OS << Dimension;
}
OS << '>';
}
Expand Down Expand Up @@ -7512,11 +7518,13 @@ void FreeQueriesInfo::printImmediateText(llvm::raw_ostream &OS, const Node *S,
return Info->printImmediateText(OS, S->getBeginLoc(), K);
}

#ifdef DPCT_DEBUG_BUILD
llvm::errs() << "Can not get FreeQueriesInfo for this FunctionDecl\n";
assert(0);
#endif // DPCT_DEBUG_BUILD

auto DFI = DeviceFunctionDecl::LinkRedecls(FD);
if (!DFI)
return;
auto Index = DpctGlobalInfo::getCudaKernelDimDFIIndexThenInc();
DpctGlobalInfo::insertCudaKernelDimDFIMap(Index, DFI);
printFreeQueriesFunctionName<std::string>(
OS, K, "{{NEEDREPLACEG" + std::to_string(Index) + "}}");
} else {
if (auto DFI = DeviceFunctionDecl::LinkRedecls(FD))
DFI->setItem();
Expand Down Expand Up @@ -7546,6 +7554,7 @@ void FreeQueriesInfo::printImmediateText(llvm::raw_ostream &OS,
(*Iter)->Infos.push_back(Idx);
Index = Iter - MacroInfos.begin();
} else {
IsMacro = false;
auto SLocInfo = DpctGlobalInfo::getLocInfo(SL);
if (SLocInfo.first != FilePath)
return;
Expand Down Expand Up @@ -7576,7 +7585,7 @@ void FreeQueriesInfo::emplaceExtraDecl() {
auto &KindNames =
getNames(static_cast<FreeQueriesKind>(FreeQueriesKind::NdItem));
OS << "auto " << KindNames.ExtraVariableName << " = ";
printFreeQueriesFunctionName(
printFreeQueriesFunctionName<unsigned>(
OS, static_cast<FreeQueriesKind>(FreeQueriesKind::NdItem), Dimension);
OS << ';' << NL << Indent;
}
Expand All @@ -7590,28 +7599,29 @@ std::string FreeQueriesInfo::getReplaceString(unsigned Num) {
bool IsMacro = isMacro(Num);
if (IsMacro) {
if (Index < MacroInfos.size()) {
return buildStringFromPrinter(printFreeQueriesFunctionName, Kind,
MacroInfos[Index]->Dimension);
return buildStringFromPrinter(printFreeQueriesFunctionName<unsigned>,
Kind, MacroInfos[Index]->Dimension);
}
#ifdef DPCT_DEBUG_BUILD
llvm::errs() << "FreeQueriesInfo index[" << Index
<< "]is larger than list size[" << InfoList.size() << "]\n";
<< "] is larger than list size[" << MacroInfos.size() << "]\n";
assert(0);
#endif // DPCT_DEBUG_BUILD
}
if (Index < InfoList.size())
return InfoList[Index]->getReplaceString(getKind(Num));
#ifdef DPCT_DEBUG_BUILD
llvm::errs() << "FreeQueriesInfo index[" << Index
<< "]is larger than list size[" << InfoList.size() << "]\n";
<< "] is larger than list size[" << InfoList.size() << "]\n";
assert(0);
#endif // DPCT_DEBUG_BUILD
return "";
}

std::string FreeQueriesInfo::getReplaceString(FreeQueriesKind K) {
if (K != FreeQueriesKind::NdItem || Counter[K] < 2)
return buildStringFromPrinter(printFreeQueriesFunctionName, K, Dimension);
return buildStringFromPrinter(printFreeQueriesFunctionName<unsigned>, K,
Dimension);
else
return getNames(K).ExtraVariableName;
}
Expand Down
2 changes: 0 additions & 2 deletions clang/lib/DPCT/RuleInfra/CallExprRewriterCommon.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -149,8 +149,6 @@ template <class SubExprT> class CastIfNotSameExprPrinter {
clang::QualType ArgType = InputArg->getType().getCanonicalType();
ArgType.removeLocalFastQualifiers(clang::Qualifiers::CVRMask);
bool NeedParen = false;
std::cout << "Arg type: " << ArgType.getAsString() << "\n";
std::cout << "Given type " << TypeInfo << "\n";
if (ArgType.getAsString() != TypeInfo) {
NeedParen = needExtraParens(SubExpr);
Stream << "(" << TypeInfo << ")";
Expand Down
24 changes: 12 additions & 12 deletions clang/test/dpct/builtin_warpSize.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: dpct --no-dpcpp-extensions=free-function-queries --format-range=none -out-root %T/builtin_warpSize %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only
// RUN: dpct --format-range=none -out-root %T/builtin_warpSize %s --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only
// RUN: FileCheck --input-file %T/builtin_warpSize/builtin_warpSize.dp.cpp --match-full-lines %s
// RUN: %if build_lit %{icpx -c -fsycl %T/builtin_warpSize/builtin_warpSize.dp.cpp -o %T/builtin_warpSize/builtin_warpSize.dp.o %}

Expand All @@ -8,7 +8,7 @@


__global__ void foo(){
// CHECK: int a = item_ct1.get_sub_group().get_local_range().get(0);
// CHECK: int a = sycl::ext::oneapi::this_work_item::get_sub_group().get_local_range().get(0);
// CHECK-NEXT: int warpSize = 1;
// CHECK-NEXT: warpSize = 2;
// CHECK-NEXT: int c= warpSize;
Expand All @@ -18,8 +18,8 @@ __global__ void foo(){
int c= warpSize;
}

// CHECK: void bar(const sycl::nd_item<3> &item_ct1){
// CHECK-NEXT: int a = sycl::max((int)item_ct1.get_sub_group().get_local_range().get(0), 0);
// CHECK: void bar(){
// CHECK-NEXT: int a = sycl::max((int)sycl::ext::oneapi::this_work_item::get_sub_group().get_local_range().get(0), 0);
// CHECK-NEXT: int warpSize = 1;
// CHECK-NEXT: int b = sycl::max(warpSize, 0);
// CHECK-NEXT: }
Expand All @@ -29,8 +29,8 @@ __global__ void bar(){
int b = max(warpSize, 0);
}

// CHECK: int tensorPos(const int ct, const sycl::nd_item<3> &item_ct1, int numLane = 0) {
// CHECK-NEXT: if (!numLane) numLane = item_ct1.get_sub_group().get_local_range().get(0);
// CHECK: int tensorPos(const int ct, int numLane = 0) {
// CHECK-NEXT: if (!numLane) numLane = sycl::ext::oneapi::this_work_item::get_sub_group().get_local_range().get(0);
// CHECK-NEXT: int r = ct * numLane;
// CHECK-NEXT: return r;
// CHECK-NEXT: }
Expand All @@ -39,18 +39,18 @@ __device__ int tensorPos(const int ct, const int numLane = warpSize) {
return r;
}

// CHECK: int tensorPos(const int ct, const sycl::nd_item<3> &item_ct1, int numLane);
// CHECK: int tensorPos(const int ct, int numLane);
__device__ int tensorPos(const int ct, const int numLane);





// CHECK: int tensorPos2(const int ct, const sycl::nd_item<3> &item_ct1, int numLane);
// CHECK: int tensorPos2(const int ct, int numLane);
__device__ int tensorPos2(const int ct, const int numLane);

// CHECK: int tensorPos2(const int ct, const sycl::nd_item<3> &item_ct1, int numLane) {
// CHECK-NEXT: if (!numLane) numLane = item_ct1.get_sub_group().get_local_range().get(0);
// CHECK: int tensorPos2(const int ct, int numLane) {
// CHECK-NEXT: if (!numLane) numLane = sycl::ext::oneapi::this_work_item::get_sub_group().get_local_range().get(0);
// CHECK-NEXT: int r = ct * numLane;
// CHECK-NEXT: return r;
// CHECK-NEXT: }
Expand All @@ -59,9 +59,9 @@ __device__ int tensorPos2(const int ct, const int numLane) {
return r;
}

// CHECK: int tensorPos2(const int ct, const sycl::nd_item<3> &item_ct1, int numLane = 0);
// CHECK: int tensorPos2(const int ct, int numLane = 0);
__device__ int tensorPos2(const int ct, const int numLane = warpSize);


// CHECK: int tensorPos3(const int ct, const sycl::nd_item<3> &item_ct1, int numLane = 0) {}
// CHECK: int tensorPos3(const int ct, int numLane = 0) {}
__device__ int tensorPos3(const int ct, const int numLane = warpSize) {}
42 changes: 25 additions & 17 deletions clang/test/dpct/macro_test.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,7 +5,7 @@
// RUN: cd %T
// RUN: rm -rf %T/macro_test_output
// RUN: mkdir %T/macro_test_output
// RUN: dpct --no-dpcpp-extensions=free-function-queries -out-root %T/macro_test_output macro_test.cu --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only
// RUN: dpct -out-root %T/macro_test_output macro_test.cu --cuda-include-path="%cuda-path/include" -- -x cuda --cuda-host-only
// RUN: FileCheck --input-file %T/macro_test_output/macro_test.dp.cpp --match-full-lines macro_test.cu
// RUN: %if build_lit %{icpx -c -fsycl -DNO_BUILD_TEST %T/macro_test_output/macro_test.dp.cpp -o %T/macro_test_output/macro_test.dp.o %}
// RUN: FileCheck --input-file %T/macro_test_output/macro_test.h --match-full-lines macro_test.h
Expand Down Expand Up @@ -67,8 +67,8 @@ __global__ void foo_kernel() {}
//CHECK-NEXT: #ifdef MACRO_CC
//CHECK-NEXT: , int c
//CHECK-NEXT: #endif
//CHECK-NEXT: , const sycl::nd_item<3> &item_ct1) {
//CHECK-NEXT: int x = item_ct1.get_group(2);
//CHECK-NEXT: ) {
//CHECK-NEXT: int x = sycl::ext::oneapi::this_work_item::get_nd_item<3>().get_group(2);
//CHECK-NEXT: }
__global__ void foo_kernel2(int a, int b
#ifdef MACRO_CC
Expand Down Expand Up @@ -288,7 +288,7 @@ int b;
//CHECK-NEXT: sycl::nd_range<3>(sycl::range<3>(1, 1, 2) * sycl::range<3>(1, 1, 2),
//CHECK-NEXT: sycl::range<3>(1, 1, 2)),
//CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) {
//CHECK-NEXT: foo_kernel2(3, 3, item_ct1);
//CHECK-NEXT: foo_kernel2(3, 3);
//CHECK-NEXT: });
foo_kernel2<<<2, 2, 0>>>(3,3
#ifdef MACRO_CC
Expand Down Expand Up @@ -435,13 +435,14 @@ FFF

}

// CHECK: #define FFFFF(aaa,bbb) void foo4(const int * __restrict__ aaa, const float * __restrict__ bbb, int *c, BBB, const sycl::nd_item<3> &item_ct1, float *sp_lj, float *sp_coul, int *ljd, double la[8][1])
// CHECK: #define FFFFF(aaa,bbb) void foo4(const int * __restrict__ aaa, const float * __restrict__ bbb, int *c, BBB, float *sp_lj, float *sp_coul, int *ljd, double la[8][1])
#define FFFFF(aaa,bbb) __device__ void foo4(const int * __restrict__ aaa, const float * __restrict__ bbb, int *c, BBB)

// CHECK: FFFFF(pos, q)
// CHECK-NEXT: {
// CHECK-EMPTY:
// CHECK-NEXT: const int tid = item_ct1.get_local_id(2);
// CHECK-NEXT: const int tid =
// CHECK-NEXT: sycl::ext::oneapi::this_work_item::get_nd_item<3>().get_local_id(2);
// CHECK-NEXT: }
FFFFF(pos, q)
{
Expand All @@ -452,13 +453,14 @@ FFFFF(pos, q)
const int tid = threadIdx.x;
}

// CHECK: #define FFFFFF(aaa,bbb) void foo5(const int * __restrict__ aaa, const float * __restrict__ bbb, const sycl::nd_item<3> &item_ct1, float *sp_lj, float *sp_coul, int *ljd, double la[8][1])
// CHECK: #define FFFFFF(aaa,bbb) void foo5(const int * __restrict__ aaa, const float * __restrict__ bbb, float *sp_lj, float *sp_coul, int *ljd, double la[8][1])
#define FFFFFF(aaa,bbb) __device__ void foo5(const int * __restrict__ aaa, const float * __restrict__ bbb)

// CHECK: FFFFFF(pos, q)
// CHECK-NEXT: {
// CHECK-EMPTY:
// CHECK-NEXT: const int tid = item_ct1.get_local_id(2);
// CHECK-NEXT: const int tid =
// CHECK-NEXT: sycl::ext::oneapi::this_work_item::get_nd_item<3>().get_local_id(2);
// CHECK-NEXT: }
FFFFFF(pos, q)
{
Expand All @@ -483,9 +485,13 @@ __device__ void foo6(AAA, BBB)

//CHECK: #define MM __umul24
//CHECK-NEXT: #define MUL(a, b) sycl::mul24((unsigned int)a, (unsigned int)b)
//CHECK-NEXT: void foo7(const sycl::nd_item<3> &item_ct1) {
//CHECK-NEXT: unsigned int tid = MUL(item_ct1.get_local_range(2), item_ct1.get_group(2)) +
//CHECK-NEXT: item_ct1.get_local_range(2);
//CHECK-NEXT: void foo7() {
//CHECK-NEXT: auto item_ct1 = sycl::ext::oneapi::this_work_item::get_nd_item<3>();
//CHECK-NEXT: unsigned int tid =
//CHECK-NEXT: MUL(sycl::ext::oneapi::this_work_item::get_nd_item<3>().get_local_range(
//CHECK-NEXT: 2),
//CHECK-NEXT: sycl::ext::oneapi::this_work_item::get_nd_item<3>().get_group(2)) +
//CHECK-NEXT: item_ct1.get_local_range(2);
//CHECK-NEXT: unsigned int tid2 = sycl::mul24((unsigned int)item_ct1.get_local_range(2),
//CHECK-NEXT: (unsigned int)item_ct1.get_group_range(2));
//CHECK-NEXT: }
Expand Down Expand Up @@ -573,7 +579,7 @@ void templatefoo2(){
CALL_KERNEL2(8, AAA)
}

//CHECK: void foo11(const sycl::nd_item<3> &item_ct1){
//CHECK: void foo11(){
//CHECK-NEXT: sycl::exp((double)(THREAD_IDX_X));
//CHECK-NEXT: }
__global__ void foo11(){
Expand Down Expand Up @@ -915,13 +921,14 @@ void foo20() {
}

//CHECK: #define CALLSHFLSYNC(x) \
//CHECK-NEXT: dpct::select_from_sub_group(item_ct1.get_sub_group(), x, 3 ^ 1);
//CHECK-NEXT: dpct::select_from_sub_group( \
//CHECK-NEXT: sycl::ext::oneapi::this_work_item::get_sub_group(), x, 3 ^ 1);
#define CALLSHFLSYNC(x) __shfl_sync(0xffffffff, x, 3 ^ 1);
//CHECK: #define CALLANYSYNC(x) \
//CHECK-NEXT: sycl::any_of_group( \
//CHECK-NEXT: item_ct1.get_sub_group(), \
//CHECK-NEXT: (0xffffffff & \
//CHECK-NEXT: (0x1 << item_ct1.get_sub_group().get_local_linear_id())) && \
//CHECK-NEXT: sycl::ext::oneapi::this_work_item::get_sub_group(), \
//CHECK-NEXT: (0xffffffff & (0x1 << sycl::ext::oneapi::this_work_item::get_sub_group() \
//CHECK-NEXT: .get_local_linear_id())) && \
//CHECK-NEXT: x != 0.0f);
#define CALLANYSYNC(x) __any_sync(0xffffffff, x != 0.0f);

Expand Down Expand Up @@ -964,7 +971,8 @@ foo23(void)
}

//CHECK: #define SHFL(x, y, z) \
//CHECK-NEXT: dpct::select_from_sub_group(item_ct1.get_sub_group(), (x), (y), (z))
//CHECK-NEXT: dpct::select_from_sub_group( \
//CHECK-NEXT: sycl::ext::oneapi::this_work_item::get_sub_group(), (x), (y), (z))
#define SHFL(x, y, z) __shfl((x), (y), (z))
__global__ void foo24(){
int i;
Expand Down
3 changes: 2 additions & 1 deletion clang/test/dpct/macro_test.h
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,5 @@
//CHECK: #define THREAD_IDX_X item_ct1.get_local_id(2)
//CHECK: #define THREAD_IDX_X \
//CHECK-NEXT: sycl::ext::oneapi::this_work_item::get_nd_item<3>().get_local_id(2)
#define THREAD_IDX_X threadIdx.x

#define STRINGIFY_(...) #__VA_ARGS__
Expand Down
Loading