Skip to content
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

Jump to bottom

[SYCLomatic] Fix the dump function with type, when type usage and type definition in the same context. #2416

Open
wants to merge 2 commits into
base: SYCLomatic
Choose a base branch
from
+361 −9
Open

[SYCLomatic] Fix the dump function with type, when type usage and type definition in the same context. #2416

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
8b7fefa
[SYCLomatic] Fix the cdump function with type, when type usage and ty…
ShengchenJ Oct 16, 2024
a347587
update lit test
ShengchenJ Oct 17, 2024
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
38 changes: 29 additions & 9 deletions clang/lib/DPCT/GenCodePinHeader.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -296,7 +296,33 @@ void GenCodePinHeaderRule::collectInfoForCodePinDumpFunction(QualType T) {
!isTypeInAnalysisScope(QT.getTypePtrOrNull())) {
continue;
}
std::string Scope = "";
std::string TypeName = QT.getAsString(PrintPolicy);
std::string TypeNameWithoutScope = "";
// For typedef type, use the alias name as the TypeName.
if (QT->isTypedefNameType()) {
TypeNameWithoutScope =
QT->getAs<TypedefType>()->getDecl()->getName().str();
} else {
PrintPolicy.SuppressScope = true;
TypeNameWithoutScope = QT.getAsString(PrintPolicy);
PrintPolicy.SuppressScope = false;
}

// When type definition and type usage are both in the namespace, the
// namespace scope can omit. Call the desugared type to get the namespace
// scope.
if (QT->getAsRecordDecl()->getDeclContext()->isNamespace()) {
TypeName = TypeNameWithoutScope;
std::string DesuagerType =
QT.getDesugaredType(Ctx).getAsString(PrintPolicy);
size_t Pos = DesuagerType.rfind("::");
if (Pos != std::string::npos) {
Scope = DesuagerType.substr(0, Pos + 2);
}
}
TypeName = Scope + TypeName;

auto &Vec = DpctGlobalInfo::getCodePinTypeInfoVec();
auto Iter =
std::find_if(Vec.begin(), Vec.end(), [&TypeName](const auto &pair) {
Expand All @@ -312,18 +338,12 @@ void GenCodePinHeaderRule::collectInfoForCodePinDumpFunction(QualType T) {
}
VarInfo.VarName = TypeName;
collectMemberInfo(QT, VarInfo, NextTypeQueue, false, PrintPolicy);
PrintPolicy.SuppressScope = 1;
std::string TypenameWithoutScope = QT.getAsString(PrintPolicy);
if (QT->isTypedefNameType())
TypenameWithoutScope =
QT->getAs<TypedefType>()->getDecl()->getName().str();
auto Pos = TypenameWithoutScope.find('<');
auto Pos = TypeNameWithoutScope.find('<');
VarInfo.VarNameWithoutScopeAndTemplateArgs =
TypenameWithoutScope.substr(0, Pos);
TypeNameWithoutScope.substr(0, Pos);
if (Pos != std::string::npos) {
VarInfo.TemplateInstArgs = TypenameWithoutScope.substr(Pos);
VarInfo.TemplateInstArgs = TypeNameWithoutScope.substr(Pos);
}
PrintPolicy.SuppressScope = 0;
Vec.push_back({TypeName, VarInfo});
}
CurrentTypeQueue = NextTypeQueue;
Expand Down
117 changes: 117 additions & 0 deletions ...st/dpct/debug_test/struct_kernel_call_in_same_namespace/codepin_autogen_util.hpp.cuda.ref
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,117 @@
//CHECK: #ifndef __DPCT_CODEPIN_AUTOGEN_UTIL__
//CHECK: #define __DPCT_CODEPIN_AUTOGEN_UTIL__

//CHECK: namespace nnn{
//CHECK: struct PP2;
//CHECK: }

//CHECK: namespace test_codepin{
//CHECK: struct P2;
//CHECK: using Point2D = P2;
//CHECK: }


//CHECK: namespace dpct {
//CHECK: namespace experimental {
//CHECK: namespace codepin {
//CHECK: namespace detail {

//CHECK: struct PP2_codepin {
//CHECK: public:
//CHECK: int x;
//CHECK: int y;
//CHECK: };
//CHECK: struct Point2D_codepin {
//CHECK: public:
//CHECK: int x;
//CHECK: int y;
//CHECK: };

//CHECK: template <> class data_ser<Point2D_codepin> {
//CHECK: public:
//CHECK: static void dump(dpctexp::codepin::detail::json_stringstream &ss, Point2D_codepin &value,
//CHECK: dpctexp::codepin::queue_t queue) {
//CHECK: auto arr = ss.array();
//CHECK: {
//CHECK: auto obj0 = arr.object();
//CHECK: obj0.key("x");
//CHECK: auto value0 = obj0.value<dpctexp::codepin::detail::json_stringstream::json_obj>();
//CHECK: dpctexp::codepin::detail::data_ser<int>::print_type_name(value0);
//CHECK: obj0.key("Data");
//CHECK: dpctexp::codepin::detail::data_ser<int>::dump(ss, value.x, queue);
//CHECK: }
//CHECK: {
//CHECK: auto obj1 = arr.object();
//CHECK: obj1.key("y");
//CHECK: auto value1 = obj1.value<dpctexp::codepin::detail::json_stringstream::json_obj>();
//CHECK: dpctexp::codepin::detail::data_ser<int>::print_type_name(value1);
//CHECK: obj1.key("Data");
//CHECK: dpctexp::codepin::detail::data_ser<int>::dump(ss, value.y, queue);
//CHECK: }
//CHECK: }
//CHECK: static void print_type_name(json_stringstream::json_obj &obj) {
//CHECK: obj.key("Type");
//CHECK: obj.value("Point2D_codepin");
//CHECK: }
//CHECK: };

//CHECK: template <> class data_ser<test_codepin::Point2D> {
//CHECK: public:
//CHECK: static void dump(dpctexp::codepin::detail::json_stringstream &ss, test_codepin::Point2D &value,
//CHECK: dpctexp::codepin::queue_t queue) {
//CHECK: Point2D_codepin& temp = reinterpret_cast<Point2D_codepin&>(value);
//CHECK: dpctexp::codepin::detail::data_ser<Point2D_codepin>::dump(ss, temp, queue);
//CHECK: }
//CHECK: static void print_type_name(json_stringstream::json_obj &obj) {
//CHECK: obj.key("Type");
//CHECK: obj.value("test_codepin::Point2D");
//CHECK: }
//CHECK: };


//CHECK: template <> class data_ser<PP2_codepin> {
//CHECK: public:
//CHECK: static void dump(dpctexp::codepin::detail::json_stringstream &ss, PP2_codepin &value,
//CHECK: dpctexp::codepin::queue_t queue) {
//CHECK: auto arr = ss.array();
//CHECK: {
//CHECK: auto obj0 = arr.object();
//CHECK: obj0.key("x");
//CHECK: auto value0 = obj0.value<dpctexp::codepin::detail::json_stringstream::json_obj>();
//CHECK: dpctexp::codepin::detail::data_ser<int>::print_type_name(value0);
//CHECK: obj0.key("Data");
//CHECK: dpctexp::codepin::detail::data_ser<int>::dump(ss, value.x, queue);
//CHECK: }
//CHECK: {
//CHECK: auto obj1 = arr.object();
//CHECK: obj1.key("y");
//CHECK: auto value1 = obj1.value<dpctexp::codepin::detail::json_stringstream::json_obj>();
//CHECK: dpctexp::codepin::detail::data_ser<int>::print_type_name(value1);
//CHECK: obj1.key("Data");
//CHECK: dpctexp::codepin::detail::data_ser<int>::dump(ss, value.y, queue);
//CHECK: }
//CHECK: }
//CHECK: static void print_type_name(json_stringstream::json_obj &obj) {
//CHECK: obj.key("Type");
//CHECK: obj.value("PP2_codepin");
//CHECK: }
//CHECK: };

//CHECK: template <> class data_ser<nnn::PP2> {
//CHECK: public:
//CHECK: static void dump(dpctexp::codepin::detail::json_stringstream &ss, nnn::PP2 &value,
//CHECK: dpctexp::codepin::queue_t queue) {
//CHECK: PP2_codepin& temp = reinterpret_cast<PP2_codepin&>(value);
//CHECK: dpctexp::codepin::detail::data_ser<PP2_codepin>::dump(ss, temp, queue);
//CHECK: }
//CHECK: static void print_type_name(json_stringstream::json_obj &obj) {
//CHECK: obj.key("Type");
//CHECK: obj.value("nnn::PP2");
//CHECK: }
//CHECK: };

//CHECK: }
//CHECK: }
//CHECK: }
//CHECK: }
//CHECK: #endif
120 changes: 120 additions & 0 deletions clang/test/dpct/debug_test/struct_kernel_call_in_same_namespace/codepin_autogen_util.hpp.ref
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,120 @@
//CHECK: #ifndef __DPCT_CODEPIN_AUTOGEN_UTIL__
//CHECK: #define __DPCT_CODEPIN_AUTOGEN_UTIL__

//CHECK: namespace nnn{
//CHECK: struct PP2;
//CHECK: }

//CHECK: namespace test_codepin{
//CHECK: struct P2;
//CHECK: using Point2D = P2;
//CHECK: }


//CHECK: namespace dpct {
//CHECK: namespace experimental {
//CHECK: namespace codepin {
//CHECK: namespace detail {

//CHECK: struct PP2_codepin {
//CHECK: public:
//CHECK: int x;
//CHECK: int y;
//CHECK: };

//CHECK: struct Point2D_codepin {
//CHECK: public:
//CHECK: int x;
//CHECK: int y;
//CHECK: };



//CHECK: template <> class data_ser<Point2D_codepin> {
//CHECK: public:
//CHECK: static void dump(dpctexp::codepin::detail::json_stringstream &ss, Point2D_codepin &value,
//CHECK: dpctexp::codepin::queue_t queue) {
//CHECK: auto arr = ss.array();
//CHECK: {
//CHECK: auto obj0 = arr.object();
//CHECK: obj0.key("x");
//CHECK: auto value0 = obj0.value<dpctexp::codepin::detail::json_stringstream::json_obj>();
//CHECK: dpctexp::codepin::detail::data_ser<int>::print_type_name(value0);
//CHECK: obj0.key("Data");
//CHECK: dpctexp::codepin::detail::data_ser<int>::dump(ss, value.x, queue);
//CHECK: }
//CHECK: {
//CHECK: auto obj1 = arr.object();
//CHECK: obj1.key("y");
//CHECK: auto value1 = obj1.value<dpctexp::codepin::detail::json_stringstream::json_obj>();
//CHECK: dpctexp::codepin::detail::data_ser<int>::print_type_name(value1);
//CHECK: obj1.key("Data");
//CHECK: dpctexp::codepin::detail::data_ser<int>::dump(ss, value.y, queue);
//CHECK: }

//CHECK: }
//CHECK: static void print_type_name(json_stringstream::json_obj &obj) {
//CHECK: obj.key("Type");
//CHECK: obj.value("Point2D_codepin");
//CHECK: }
//CHECK: };

//CHECK: template <> class data_ser<test_codepin::Point2D> {
//CHECK: public:
//CHECK: static void dump(dpctexp::codepin::detail::json_stringstream &ss, test_codepin::Point2D &value,
//CHECK: dpctexp::codepin::queue_t queue) {
//CHECK: Point2D_codepin& temp = reinterpret_cast<Point2D_codepin&>(value);
//CHECK: dpctexp::codepin::detail::data_ser<Point2D_codepin>::dump(ss, temp, queue);
//CHECK: }
//CHECK: static void print_type_name(json_stringstream::json_obj &obj) {
//CHECK: obj.key("Type");
//CHECK: obj.value("test_codepin::Point2D");
//CHECK: }
//CHECK: };


//CHECK: template <> class data_ser<PP2_codepin> {
//CHECK: public:
//CHECK: static void dump(dpctexp::codepin::detail::json_stringstream &ss, PP2_codepin &value,
//CHECK: dpctexp::codepin::queue_t queue) {
//CHECK: auto arr = ss.array();
//CHECK: {
//CHECK: auto obj0 = arr.object();
//CHECK: obj0.key("x");
//CHECK: auto value0 = obj0.value<dpctexp::codepin::detail::json_stringstream::json_obj>();
//CHECK: dpctexp::codepin::detail::data_ser<int>::print_type_name(value0);
//CHECK: obj0.key("Data");
//CHECK: dpctexp::codepin::detail::data_ser<int>::dump(ss, value.x, queue);
//CHECK: }
//CHECK: {
//CHECK: auto obj1 = arr.object();
//CHECK: obj1.key("y");
//CHECK: auto value1 = obj1.value<dpctexp::codepin::detail::json_stringstream::json_obj>();
//CHECK: dpctexp::codepin::detail::data_ser<int>::print_type_name(value1);
//CHECK: obj1.key("Data");
//CHECK: dpctexp::codepin::detail::data_ser<int>::dump(ss, value.y, queue);
//CHECK: }
//CHECK: }
//CHECK: static void print_type_name(json_stringstream::json_obj &obj) {
//CHECK: obj.key("Type");
//CHECK: obj.value("PP2_codepin");
//CHECK: }
//CHECK: };
//CHECK: template <> class data_ser<nnn::PP2> {
//CHECK: public:
//CHECK: static void dump(dpctexp::codepin::detail::json_stringstream &ss, nnn::PP2 &value,
//CHECK: dpctexp::codepin::queue_t queue) {
//CHECK: PP2_codepin& temp = reinterpret_cast<PP2_codepin&>(value);
//CHECK: dpctexp::codepin::detail::data_ser<PP2_codepin>::dump(ss, temp, queue);
//CHECK: }
//CHECK: static void print_type_name(json_stringstream::json_obj &obj) {
//CHECK: obj.key("Type");
//CHECK: obj.value("nnn::PP2");
//CHECK: }
//CHECK: };

//CHECK: }
//CHECK: }
//CHECK: }
//CHECK: }
//CHECK: #endif
95 changes: 95 additions & 0 deletions clang/test/dpct/debug_test/struct_kernel_call_in_same_namespace/test.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,95 @@
// RUN: dpct --format-range=none --enable-codepin -out-root %T/debug_test/struct_kernel_call_in_same_namespace %s --cuda-include-path="%cuda-path/include" -- -std=c++17 -x cuda --cuda-host-only
// RUN: FileCheck %S/codepin_autogen_util.hpp.ref --match-full-lines --input-file %T/debug_test/struct_kernel_call_in_same_namespace_codepin_sycl/codepin_autogen_util.hpp
// RUN: FileCheck %S/codepin_autogen_util.hpp.cuda.ref --match-full-lines --input-file %T/debug_test/struct_kernel_call_in_same_namespace_codepin_cuda/codepin_autogen_util.hpp
// RUN: FileCheck %s --match-full-lines --input-file %T/debug_test/struct_kernel_call_in_same_namespace_codepin_sycl/test.dp.cpp
// RUN: FileCheck %s --match-full-lines --input-file %T/debug_test/struct_kernel_call_in_same_namespace_codepin_cuda/test.cu
// RUN: %if build_lit %{icpx -c -fsycl %T/debug_test/struct_kernel_call_in_same_namespace_codepin_sycl/test.dp.cpp -o %T/debug_test/struct_kernel_call_in_same_namespace_codepin_sycl/test.dp.o %}
#include <cuda.h>
#include <iostream>
namespace test_codepin {
struct P2 {
int x;
int y;
};
using Point2D = P2;
};

//CHECK: namespace nnn {
namespace nnn {
struct PP2{
int x;
int y;
} ;

using INT = int;

__global__ void kernel2d(test_codepin::Point2D *a, test_codepin::Point2D *b, test_codepin::Point2D *c) {
int i = threadIdx.x;
c[i].x = a[i].x + b[i].x;
c[i].y = a[i].y + b[i].y;
}

__global__ void kernel2d_2(PP2 *a, PP2 *b, PP2 *c) {
int i = threadIdx.x;
c[i].x = a[i].x + b[i].x;
c[i].y = a[i].y + b[i].y;
}

__global__ void kerneel_int(INT *a, INT *b, INT *c) {
int i = threadIdx.x;
c[i] = a[i] + b[i];
}

#define NUM 10
void test() {
test_codepin::Point2D h_2d[NUM];
for (int i = 0; i < NUM; i++) {
h_2d[i].x = i;
h_2d[i].y = i;
}
test_codepin::Point2D *d_a2d, *d_b2d, *d_c2d;
cudaMalloc(&d_a2d, sizeof(test_codepin::Point2D) * NUM);
cudaMalloc(&d_b2d, sizeof(test_codepin::Point2D) * NUM);
cudaMalloc(&d_c2d, sizeof(test_codepin::Point2D) * NUM);
cudaMemcpy(d_a2d, h_2d, sizeof(test_codepin::Point2D) * NUM, cudaMemcpyHostToDevice);
cudaMemcpy(d_b2d, h_2d, sizeof(test_codepin::Point2D) * NUM, cudaMemcpyHostToDevice);
cudaMemcpy(d_c2d, h_2d, sizeof(test_codepin::Point2D) * NUM, cudaMemcpyHostToDevice);
kernel2d<<<1, NUM>>>(d_a2d, d_b2d, d_c2d);
cudaDeviceSynchronize();

PP2 h_pp2_2d[NUM];
for (int i = 0; i < NUM; i++) {
h_pp2_2d[i].x = i;
h_pp2_2d[i].y = i;
}
PP2 *d_pp2_a2d, *d_pp2_b2d, *d_pp2_c2d;
cudaMalloc(&d_pp2_a2d, sizeof(PP2) * NUM);
cudaMalloc(&d_pp2_b2d, sizeof(PP2) * NUM);
cudaMalloc(&d_pp2_c2d, sizeof(PP2) * NUM);
cudaMemcpy(d_pp2_a2d, h_pp2_2d, sizeof(PP2) * NUM, cudaMemcpyHostToDevice);
cudaMemcpy(d_pp2_b2d, h_pp2_2d, sizeof(PP2) * NUM, cudaMemcpyHostToDevice);
cudaMemcpy(d_pp2_c2d, h_pp2_2d, sizeof(PP2) * NUM, cudaMemcpyHostToDevice);
kernel2d_2<<<1, NUM>>>(d_pp2_a2d, d_pp2_b2d, d_pp2_c2d);
cudaDeviceSynchronize();

INT h_int_2d[NUM];
for (int i = 0; i < NUM; i++) {
h_int_2d[i] = i;
}
INT *d_int_a2d, *d_int_b2d, *d_int_c2d;
cudaMalloc(&d_int_a2d, sizeof(INT) * NUM);
cudaMalloc(&d_int_b2d, sizeof(INT) * NUM);
cudaMalloc(&d_int_c2d, sizeof(INT) * NUM);
cudaMemcpy(d_int_a2d, h_int_2d, sizeof(INT) * NUM, cudaMemcpyHostToDevice);
cudaMemcpy(d_int_b2d, h_int_2d, sizeof(INT) * NUM, cudaMemcpyHostToDevice);
cudaMemcpy(d_int_c2d, h_int_2d, sizeof(INT) * NUM, cudaMemcpyHostToDevice);
kerneel_int<<<1, NUM>>>(d_int_a2d, d_int_b2d, d_int_c2d);
cudaDeviceSynchronize();
}
}; // namespace test_codepin
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
}; // namespace test_codepin
}; // namespace nnn


int main() {


return 0;
}