diff --git a/clang/lib/DPCT/AnalysisInfo.cpp b/clang/lib/DPCT/AnalysisInfo.cpp index e13d7925bb54..cd456943f3a2 100644 --- a/clang/lib/DPCT/AnalysisInfo.cpp +++ b/clang/lib/DPCT/AnalysisInfo.cpp @@ -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; @@ -3298,7 +3301,7 @@ std::string MemVarInfo::getArgName() { return getName(); } const std::string MemVarInfo::ExternVariableName = "dpct_local"; -std::unordered_map MemVarInfo::AnonymousTypeDeclStmtMap; +std::unordered_map MemVarInfo::AnonymousTypeDeclStmtMap; ///// class TextureTypeInfo ///// TextureTypeInfo::TextureTypeInfo(std::string &&DataType, int TexType) { setDataTypeAndTexType(std::move(DataType), TexType); diff --git a/clang/lib/DPCT/AnalysisInfo.h b/clang/lib/DPCT/AnalysisInfo.h index b2712ab89942..920a33fe8714 100644 --- a/clang/lib/DPCT/AnalysisInfo.h +++ b/clang/lib/DPCT/AnalysisInfo.h @@ -1919,7 +1919,7 @@ class MemVarInfo : public VarInfo { const DeclStmt *DeclStmtOfVarType = nullptr; std::string LocalTypeName = ""; - static std::unordered_map AnonymousTypeDeclStmtMap; + static std::unordered_map AnonymousTypeDeclStmtMap; bool UsedBySymbolAPIFlag = false; bool UseHelperFuncFlag = true; bool UseDeviceGlobalFlag = false; diff --git a/clang/test/dpct/shared_union_in_kernel/common.cuh b/clang/test/dpct/shared_union_in_kernel/common.cuh new file mode 100644 index 000000000000..76ad4ee164fc --- /dev/null +++ b/clang/test/dpct/shared_union_in_kernel/common.cuh @@ -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; +} diff --git a/clang/test/dpct/shared_union_in_kernel/compile_commands.json b/clang/test/dpct/shared_union_in_kernel/compile_commands.json new file mode 100644 index 000000000000..f0acbfee9141 --- /dev/null +++ b/clang/test/dpct/shared_union_in_kernel/compile_commands.json @@ -0,0 +1,12 @@ +[ + { + "command": "nvcc -c main.cu", + "directory": ".", + "file": "main.cu" + }, + { + "command": "nvcc -c test.cu", + "directory": ".", + "file": "test.cu" + } +] diff --git a/clang/test/dpct/shared_union_in_kernel/main.cu b/clang/test/dpct/shared_union_in_kernel/main.cu new file mode 100644 index 000000000000..e374c27aa98a --- /dev/null +++ b/clang/test/dpct/shared_union_in_kernel/main.cu @@ -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; } diff --git a/clang/test/dpct/shared_union_in_kernel/test.cu b/clang/test/dpct/shared_union_in_kernel/test.cu new file mode 100644 index 000000000000..df1e6332a0af --- /dev/null +++ b/clang/test/dpct/shared_union_in_kernel/test.cu @@ -0,0 +1,3 @@ +// RUN: echo "empty command" + +#include "common.cuh"