Skip to content

Commit

Permalink
[SYCLomatic] Fix migration of shared union in header file.
Browse files Browse the repository at this point in the history
Signed-off-by: Tang, Jiajun [email protected]
  • Loading branch information
tangjj11 committed May 30, 2024
1 parent b2e5588 commit 4aa6dc6
Show file tree
Hide file tree
Showing 6 changed files with 37 additions and 4 deletions.
9 changes: 6 additions & 3 deletions clang/lib/DPCT/AnalysisInfo.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -2844,14 +2844,17 @@ MemVarInfo::MemVarInfo(unsigned Offset,
if (DS1 && DS2 && DS1 == DS2) {
IsAnonymousType = true;
DeclStmtOfVarType = DS2;
auto Iter = AnonymousTypeDeclStmtMap.find(DS2);
const auto LocInfo = DpctGlobalInfo::getLocInfo(DS2->getBeginLoc());
const auto LocStr = LocInfo.first.getCanonicalPath().str() +
std::to_string(LocInfo.second);
auto Iter = AnonymousTypeDeclStmtMap.find(LocStr);
if (Iter != AnonymousTypeDeclStmtMap.end()) {
LocalTypeName = "type_ct" + std::to_string(Iter->second);
} else {
LocalTypeName =
"type_ct" + std::to_string(AnonymousTypeDeclStmtMap.size() + 1);
AnonymousTypeDeclStmtMap.insert(
std::make_pair(DS2, AnonymousTypeDeclStmtMap.size() + 1));
std::make_pair(LocStr, AnonymousTypeDeclStmtMap.size() + 1));
}
} else if (DS2) {
DeclStmtOfVarType = DS2;
Expand Down Expand Up @@ -3298,7 +3301,7 @@ std::string MemVarInfo::getArgName() {
return getName();
}
const std::string MemVarInfo::ExternVariableName = "dpct_local";
std::unordered_map<const DeclStmt *, int> MemVarInfo::AnonymousTypeDeclStmtMap;
std::unordered_map<std::string, int> MemVarInfo::AnonymousTypeDeclStmtMap;
///// class TextureTypeInfo /////
TextureTypeInfo::TextureTypeInfo(std::string &&DataType, int TexType) {
setDataTypeAndTexType(std::move(DataType), TexType);
Expand Down
2 changes: 1 addition & 1 deletion clang/lib/DPCT/AnalysisInfo.h
Original file line number Diff line number Diff line change
Expand Up @@ -1919,7 +1919,7 @@ class MemVarInfo : public VarInfo {
const DeclStmt *DeclStmtOfVarType = nullptr;
std::string LocalTypeName = "";

static std::unordered_map<const DeclStmt *, int> AnonymousTypeDeclStmtMap;
static std::unordered_map<std::string, int> AnonymousTypeDeclStmtMap;
bool UsedBySymbolAPIFlag = false;
bool UseHelperFuncFlag = true;
bool UseDeviceGlobalFlag = false;
Expand Down
7 changes: 7 additions & 0 deletions clang/test/dpct/shared_union_in_kernel/common.cuh
Original file line number Diff line number Diff line change
@@ -0,0 +1,7 @@
// CHECK: void f(uint8_t *shared_storage_ct1) {
// CHECK-NEXT: union type_ct1{ int x; };
// CHECK-NEXT: type_ct1 &shared_storage = *(type_ct1 *)shared_storage_ct1;
// CHECK-NEXT: }
__global__ void f() {
__shared__ union { int x; } shared_storage;
}
12 changes: 12 additions & 0 deletions clang/test/dpct/shared_union_in_kernel/compile_commands.json
Original file line number Diff line number Diff line change
@@ -0,0 +1,12 @@
[
{
"command": "nvcc -c main.cu",
"directory": ".",
"file": "main.cu"
},
{
"command": "nvcc -c test.cu",
"directory": ".",
"file": "test.cu"
}
]
8 changes: 8 additions & 0 deletions clang/test/dpct/shared_union_in_kernel/main.cu
Original file line number Diff line number Diff line change
@@ -0,0 +1,8 @@
// RUN: cp %S/* .
// RUN: dpct --format-range=none -p=. -out-root=%T --cuda-include-path="%cuda-path/include" -- -x -std=c++14
// RUN: FileCheck --input-file %T/common.dp.hpp --match-full-lines %S/common.cuh
// RUN: %if build_lit %{icpx -c -fsycl %T/main.dp.cpp %T/test.dp.cpp %}

#include "common.cuh"

int main() { return 0; }
3 changes: 3 additions & 0 deletions clang/test/dpct/shared_union_in_kernel/test.cu
Original file line number Diff line number Diff line change
@@ -0,0 +1,3 @@
// RUN: echo "empty command"

#include "common.cuh"

0 comments on commit 4aa6dc6

Please sign in to comment.