From d842dfeafbf2a2bc79edd0f613dfeea3b0af535e Mon Sep 17 00:00:00 2001 From: abhilash1910 Date: Fri, 15 Sep 2023 12:02:38 -0700 Subject: [PATCH 01/19] cub::blockload migration --- clang/lib/DPCT/CUBAPIMigration.cpp | 16 ++++++++++------ clang/lib/DPCT/ExprAnalysis.cpp | 2 +- 2 files changed, 11 insertions(+), 7 deletions(-) diff --git a/clang/lib/DPCT/CUBAPIMigration.cpp b/clang/lib/DPCT/CUBAPIMigration.cpp index 74c4a2753157..dfef83d73b80 100644 --- a/clang/lib/DPCT/CUBAPIMigration.cpp +++ b/clang/lib/DPCT/CUBAPIMigration.cpp @@ -558,14 +558,14 @@ void CubDeviceLevelRule::removeRedundantTempVar(const CallExpr *CE) { void CubRule::registerMatcher(ast_matchers::MatchFinder &MF) { MF.addMatcher( typeLoc(loc(qualType(hasDeclaration(namedDecl(hasAnyName( - "WarpScan", "WarpReduce", "BlockScan", "BlockReduce")))))) + "WarpScan", "WarpReduce", "BlockScan", "BlockReduce", "BlockLoad")))))) .bind("TypeLoc"), this); MF.addMatcher( typedefDecl( hasType(hasCanonicalType(qualType(hasDeclaration(namedDecl(hasAnyName( - "WarpScan", "WarpReduce", "BlockScan", "BlockReduce"))))))) + "WarpScan", "WarpReduce", "BlockScan", "BlockReduce", "BlockLoad"))))))) .bind("TypeDefDecl"), this); @@ -684,7 +684,8 @@ void CubRule::processCubDeclStmt(const DeclStmt *DS) { ObjTypeStr.find("class cub::WarpReduce") == 0) { Repl = DpctGlobalInfo::getSubGroup(DRE); } else if (ObjTypeStr.find("class cub::BlockScan") == 0 || - ObjTypeStr.find("class cub::BlockReduce") == 0) { + ObjTypeStr.find("class cub::BlockReduce") == 0 || + ObjTypeStr.find("class cub::BlockLoad") == 0) { Repl = DpctGlobalInfo::getGroup(DRE); } else { continue; @@ -749,7 +750,8 @@ void CubRule::processCubTypeDef(const TypedefDecl *TD) { !(ObjTypeStr.find("class cub::WarpScan") == 0 || ObjTypeStr.find("class cub::WarpReduce") == 0 || ObjTypeStr.find("class cub::BlockScan") == 0 || - ObjTypeStr.find("class cub::BlockReduce") == 0)) { + ObjTypeStr.find("class cub::BlockReduce") == 0 || + ObjTypeStr.find("class cub::BlockLoad") == 0)) { DeleteFlag = false; break; } @@ -1304,7 +1306,8 @@ void CubRule::processCubMemberCall(const CXXMemberCallExpr *MC) { ObjTypeStr.find("class cub::WarpReduce") == 0) { processWarpLevelMemberCall(MC); } else if (ObjTypeStr.find("class cub::BlockScan") == 0 || - ObjTypeStr.find("class cub::BlockReduce") == 0) { + ObjTypeStr.find("class cub::BlockReduce") == 0 || + ObjTypeStr.find("class cub::BlockLoad") == 0) { processBlockLevelMemberCall(MC); } else { report(MC->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false, ObjTypeStr); @@ -1328,7 +1331,8 @@ void CubRule::processTypeLoc(const TypeLoc *TL) { MapNames::getClNamespace() + "sub_group", SM)); } else if (TypeName.find("class cub::BlockScan") == 0 || - TypeName.find("class cub::BlockReduce") == 0) { + TypeName.find("class cub::BlockReduce") == 0 || + ObjTypeStr.find("class cub::BlockLoad") == 0) { auto DeviceFuncDecl = DpctGlobalInfo::findAncestor(TL); if (DeviceFuncDecl && (DeviceFuncDecl->hasAttr() || DeviceFuncDecl->hasAttr())) { diff --git a/clang/lib/DPCT/ExprAnalysis.cpp b/clang/lib/DPCT/ExprAnalysis.cpp index c4e0b080c42e..8e840af1efb8 100644 --- a/clang/lib/DPCT/ExprAnalysis.cpp +++ b/clang/lib/DPCT/ExprAnalysis.cpp @@ -1170,7 +1170,7 @@ void ExprAnalysis::analyzeType(TypeLoc TL, const Expr *CSCE, } } if (OS.str() != "cub::WarpScan" && OS.str() != "cub::WarpReduce" && - OS.str() != "cub::BlockReduce" && OS.str() != "cub::BlockScan") { + OS.str() != "cub::BlockReduce" && OS.str() != "cub::BlockScan" && OS.str != "cub::BlockLoad") { SR.setEnd(TSTL.getTemplateNameLoc()); } analyzeTemplateSpecializationType(TSTL); From c014557be37c53d376783983d90710f601de1ea1 Mon Sep 17 00:00:00 2001 From: abhilash1910 Date: Sun, 17 Sep 2023 22:21:43 -0700 Subject: [PATCH 02/19] add blockload + fix --- clang/lib/DPCT/APINames_CUB.inc | 2 +- clang/lib/DPCT/ExprAnalysis.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/DPCT/APINames_CUB.inc b/clang/lib/DPCT/APINames_CUB.inc index 022a84f1844f..76c90626d476 100644 --- a/clang/lib/DPCT/APINames_CUB.inc +++ b/clang/lib/DPCT/APINames_CUB.inc @@ -107,7 +107,6 @@ ENTRY_MEMBER_FUNCTION(cub::BlockExchange, cub::BlockExchange, ScatterToStripedFl ENTRY_MEMBER_FUNCTION(cub::BlockHistogram, cub::BlockHistogram, InitHistogram, InitHistogram, false, NO_FLAG, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::BlockHistogram, cub::BlockHistogram, Histogram, Histogram, false, NO_FLAG, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::BlockHistogram, cub::BlockHistogram, Composite, Composite, false, NO_FLAG, P4, "Comment") -ENTRY_MEMBER_FUNCTION(cub::BlockLoad, cub::BlockLoad, Load, Load, false, NO_FLAG, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::BlockStore, cub::BlockStore, Store, Store, false, NO_FLAG, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::BlockRadixSort, cub::BlockRadixSort, Sort, Sort, false, NO_FLAG, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::BlockRadixSort, cub::BlockRadixSort, SortDescending, SortDescending, false, NO_FLAG, P4, "Comment") @@ -123,6 +122,7 @@ ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Offset, Offset, fals ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Rotate, Rotate, false, NO_FLAG, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Up, Up, false, NO_FLAG, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Down, Down, false, NO_FLAG, P4, "Comment") +ENTRY_MEMBER_FUNCTION(cub::BlockLoad, cub::BlockLoad, Load, Load, true, NO_FLAG, P4, "Comment") // Device Level diff --git a/clang/lib/DPCT/ExprAnalysis.cpp b/clang/lib/DPCT/ExprAnalysis.cpp index 8e840af1efb8..a8f870a581f4 100644 --- a/clang/lib/DPCT/ExprAnalysis.cpp +++ b/clang/lib/DPCT/ExprAnalysis.cpp @@ -1170,7 +1170,7 @@ void ExprAnalysis::analyzeType(TypeLoc TL, const Expr *CSCE, } } if (OS.str() != "cub::WarpScan" && OS.str() != "cub::WarpReduce" && - OS.str() != "cub::BlockReduce" && OS.str() != "cub::BlockScan" && OS.str != "cub::BlockLoad") { + OS.str() != "cub::BlockReduce" && OS.str() != "cub::BlockScan" && OS.str() != "cub::BlockLoad") { SR.setEnd(TSTL.getTemplateNameLoc()); } analyzeTemplateSpecializationType(TSTL); From aadbb89e6f334c308cecee2542dc5a4fa4a09fcb Mon Sep 17 00:00:00 2001 From: abhilash1910 Date: Mon, 18 Sep 2023 01:11:19 -0700 Subject: [PATCH 03/19] add block load sample --- clang/lib/DPCT/APINames_CUB.inc | 2 +- clang/test/dpct/cub/blocklevel/blockload.cu | 83 +++++++++++++++++++++ 2 files changed, 84 insertions(+), 1 deletion(-) create mode 100644 clang/test/dpct/cub/blocklevel/blockload.cu diff --git a/clang/lib/DPCT/APINames_CUB.inc b/clang/lib/DPCT/APINames_CUB.inc index 76c90626d476..1f90d76369b5 100644 --- a/clang/lib/DPCT/APINames_CUB.inc +++ b/clang/lib/DPCT/APINames_CUB.inc @@ -122,7 +122,7 @@ ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Offset, Offset, fals ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Rotate, Rotate, false, NO_FLAG, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Up, Up, false, NO_FLAG, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Down, Down, false, NO_FLAG, P4, "Comment") -ENTRY_MEMBER_FUNCTION(cub::BlockLoad, cub::BlockLoad, Load, Load, true, NO_FLAG, P4, "Comment") +ENTRY_MEMBER_FUNCTION(cub::BlockLoad, cub::BlockLoad, Load, Load, true, NO_FLAG, P4, "Successful") // Device Level diff --git a/clang/test/dpct/cub/blocklevel/blockload.cu b/clang/test/dpct/cub/blocklevel/blockload.cu new file mode 100644 index 000000000000..2a33b8bcdde9 --- /dev/null +++ b/clang/test/dpct/cub/blocklevel/blockload.cu @@ -0,0 +1,83 @@ +// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2 +// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2 +// RUN: dpct -in-root %S -out-root %T/blocklevel/blockload %S/blockload.cu --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/blocklevel/blockscan/blockload.dp.cpp --match-full-lines %s + +#include +#include + +#include +#include + +#define WARP_SIZE 32 + +const int N = 256; +const int BlockSize = 128; +const int ItemsPerThread = 4; + + +void init_data(int* data, int num) { + for(int i = 0; i < num; i++) + data[i] = i; +} +void verify_data(int* data, int num) { + return; +} +void print_data(int* data, int num) { + for (int i = 0; i < num; i++) { + std::cout << data[i] << " "; + } + std::cout << std::endl; +} + +//CHECK: void SumKernel(int* data, +//CHECK-NEXT: const sycl::nd_item<3> &item_ct1) { +//CHECK-EMPTY: +//CHECK-NEXT: int threadid = item_ct1.get_local_id(2); +//CHECK-EMPTY: +//CHECK-NEXT: int input = data[threadid]; +//CHECK-NEXT: int output = 0; +//CHECK-NEXT: output = sycl::reduce_over_group(item_ct1.get_group(), input, sycl::plus<>()); +//CHECK-NEXT: data[threadid] = output; +//CHECK-NEXT:} + +__global__ void BlockLoadKernel(int *d_data) +{ + // Specialize BlockLoad for a 1D block of 128 threads owning 4 integer items each + typedef cub::BlockLoad BlockLoad; + // Allocate shared memory for BlockLoad + __shared__ typename BlockLoad::TempStorage temp_storage; + // Load a segment of consecutive items that are blocked across threads + int thread_data[ItemsPerThread]; + int offset = threadIdx.x * ItemsPerThread; + BlockLoad(temp_storage).Load(d_data + offset, thread_data); + + // Print loaded data + printf("Thread %d loaded: %d %d %d %d\n", threadIdx.x, thread_data[0], thread_data[1], thread_data[2], thread_data[3]); +} + +int main() +{ + int h_data[N]; + init_data(h_data, N); + int *d_data; + cudaMalloc((void**)&d_data, N * sizeof(int)); + cudaMemcpy(d_data, h_data, N * sizeof(int), cudaMemcpyHostToDevice); + + //CHECK: q_ct1.parallel_for( + //CHECK-NEXT: sycl::nd_range<3>(GridSize * BlockSize, BlockSize), + //CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { + //CHECK-NEXT: BlockLoadKernel(dev_data, item_ct1); + //CHECK-NEXT: }); + + dim3 block(BlockSize); + dim3 grid((N + BlockSize - 1) / BlockSize); + + BlockLoadKernel<<>>(d_data); + cudaDeviceSynchronize(); + //verify_data(d_data, N); + + cudaFree(d_data); + + return 0; +} \ No newline at end of file From 3e83abbc5c4df01b232a7d72c86119e62d5346e9 Mon Sep 17 00:00:00 2001 From: abhilash1910 Date: Mon, 18 Sep 2023 03:29:45 -0700 Subject: [PATCH 04/19] add line --- clang/test/dpct/cub/blocklevel/blockload.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/dpct/cub/blocklevel/blockload.cu b/clang/test/dpct/cub/blocklevel/blockload.cu index 2a33b8bcdde9..0943a6371eec 100644 --- a/clang/test/dpct/cub/blocklevel/blockload.cu +++ b/clang/test/dpct/cub/blocklevel/blockload.cu @@ -80,4 +80,4 @@ int main() cudaFree(d_data); return 0; -} \ No newline at end of file +} From c8568373e0f9f162f2bb8a9bcfb72e2ec170e594 Mon Sep 17 00:00:00 2001 From: abhilash1910 Date: Tue, 19 Sep 2023 00:16:54 -0700 Subject: [PATCH 05/19] add prototype kernel api --- clang/lib/DPCT/CUBAPIMigration.cpp | 11 +++++++++++ clang/test/dpct/cub/blocklevel/blockload.cu | 4 ++-- 2 files changed, 13 insertions(+), 2 deletions(-) diff --git a/clang/lib/DPCT/CUBAPIMigration.cpp b/clang/lib/DPCT/CUBAPIMigration.cpp index dfef83d73b80..c3306c0d489e 100644 --- a/clang/lib/DPCT/CUBAPIMigration.cpp +++ b/clang/lib/DPCT/CUBAPIMigration.cpp @@ -1141,6 +1141,17 @@ void CubRule::processBlockLevelMemberCall(const CXXMemberCallExpr *BlockMC) { CubParamAs << GroupOrWorkitem << InEA.getReplacedString() << OpRepl; Repl = NewFuncName + "(" + ParamList + ")"; emplaceTransformation(new ReplaceStmt(BlockMC, Repl)); + } else if (FuncName == "Load") { + + GroupOrWorkitem = DpctGlobalInfo::getItem(BlockMC); + NewFuncName = Mapnames::getClNamespace() + "load"; + const Expr *InData = FuncArgs[0]; + ExprAnalysis InEA(InData); + OpRepl = getOpRepl(nullptr); + CubParamAs << GroupOrWorkitem << InEA.getReplacedString() << OpRepl; + Repl = NewFuncName + "(" + ParamList + ")"; + emplaceTransformation(new ReplaceStmt(BlockMC, Repl)); + } } diff --git a/clang/test/dpct/cub/blocklevel/blockload.cu b/clang/test/dpct/cub/blocklevel/blockload.cu index 0943a6371eec..927e806e6bf3 100644 --- a/clang/test/dpct/cub/blocklevel/blockload.cu +++ b/clang/test/dpct/cub/blocklevel/blockload.cu @@ -30,14 +30,14 @@ void print_data(int* data, int num) { std::cout << std::endl; } -//CHECK: void SumKernel(int* data, +//CHECK: void BlockLoadKernel(int* data, //CHECK-NEXT: const sycl::nd_item<3> &item_ct1) { //CHECK-EMPTY: //CHECK-NEXT: int threadid = item_ct1.get_local_id(2); //CHECK-EMPTY: //CHECK-NEXT: int input = data[threadid]; //CHECK-NEXT: int output = 0; -//CHECK-NEXT: output = sycl::reduce_over_group(item_ct1.get_group(), input, sycl::plus<>()); +//CHECK-NEXT: output = sycl::load(item_ct1.get_group(), input, sycl::plus<>()); //CHECK-NEXT: data[threadid] = output; //CHECK-NEXT:} From 8b66f8ec7ef165760fd80746961be0bcec56e287 Mon Sep 17 00:00:00 2001 From: abhilash1910 Date: Tue, 19 Sep 2023 21:41:47 -0700 Subject: [PATCH 06/19] fix bugs --- clang/lib/DPCT/CUBAPIMigration.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/DPCT/CUBAPIMigration.cpp b/clang/lib/DPCT/CUBAPIMigration.cpp index c3306c0d489e..8b6530f3ff09 100644 --- a/clang/lib/DPCT/CUBAPIMigration.cpp +++ b/clang/lib/DPCT/CUBAPIMigration.cpp @@ -1144,7 +1144,7 @@ void CubRule::processBlockLevelMemberCall(const CXXMemberCallExpr *BlockMC) { } else if (FuncName == "Load") { GroupOrWorkitem = DpctGlobalInfo::getItem(BlockMC); - NewFuncName = Mapnames::getClNamespace() + "load"; + NewFuncName = MapNames::getClNamespace() + "load"; const Expr *InData = FuncArgs[0]; ExprAnalysis InEA(InData); OpRepl = getOpRepl(nullptr); @@ -1343,7 +1343,7 @@ void CubRule::processTypeLoc(const TypeLoc *TL) { SM)); } else if (TypeName.find("class cub::BlockScan") == 0 || TypeName.find("class cub::BlockReduce") == 0 || - ObjTypeStr.find("class cub::BlockLoad") == 0) { + TypeName.find("class cub::BlockLoad") == 0) { auto DeviceFuncDecl = DpctGlobalInfo::findAncestor(TL); if (DeviceFuncDecl && (DeviceFuncDecl->hasAttr() || DeviceFuncDecl->hasAttr())) { From 8b5589ac075a1ceeb0e076b9dd0be74d11e64071 Mon Sep 17 00:00:00 2001 From: abhilash1910 Date: Tue, 19 Sep 2023 21:46:08 -0700 Subject: [PATCH 07/19] fix bugs --- clang/lib/DPCT/CUBAPIMigration.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/DPCT/CUBAPIMigration.cpp b/clang/lib/DPCT/CUBAPIMigration.cpp index 8b6530f3ff09..f6bd2591060c 100644 --- a/clang/lib/DPCT/CUBAPIMigration.cpp +++ b/clang/lib/DPCT/CUBAPIMigration.cpp @@ -582,7 +582,7 @@ void CubRule::registerMatcher(ast_matchers::MatchFinder &MF) { MF.addMatcher(cxxMemberCallExpr(has(memberExpr(member(hasAnyName( "InclusiveSum", "ExclusiveSum", "InclusiveScan", "ExclusiveScan", - "Reduce", "Sum", "Broadcast", "Scan"))))) + "Reduce", "Sum", "Broadcast", "Scan", "Load"))))) .bind("MemberCall"), this); @@ -1144,7 +1144,7 @@ void CubRule::processBlockLevelMemberCall(const CXXMemberCallExpr *BlockMC) { } else if (FuncName == "Load") { GroupOrWorkitem = DpctGlobalInfo::getItem(BlockMC); - NewFuncName = MapNames::getClNamespace() + "load"; + NewFuncName = MapNames::getClNamespace() + "Load"; const Expr *InData = FuncArgs[0]; ExprAnalysis InEA(InData); OpRepl = getOpRepl(nullptr); From d0a0e8bc5e0176a6bb6b23f66c9a9c635fb62c3e Mon Sep 17 00:00:00 2001 From: Abhilash Majumder <30946547+abhilash1910@users.noreply.github.com> Date: Thu, 11 Jul 2024 11:10:47 +0530 Subject: [PATCH 08/19] modify to store --- clang/lib/DPCT/CUBAPIMigration.cpp | 19 ++++++++----------- 1 file changed, 8 insertions(+), 11 deletions(-) diff --git a/clang/lib/DPCT/CUBAPIMigration.cpp b/clang/lib/DPCT/CUBAPIMigration.cpp index 067ca823e6d2..11ad50b17238 100644 --- a/clang/lib/DPCT/CUBAPIMigration.cpp +++ b/clang/lib/DPCT/CUBAPIMigration.cpp @@ -192,7 +192,8 @@ void CubIntrinsicRule::registerMatcher(ast_matchers::MatchFinder &MF) { "WarpId", "SyncStream", "CurrentDevice", "DeviceCount", "DeviceCountUncached", "DeviceCountCachedValue", "PtxVersion", "PtxVersionUncached", "SmVersion", - "SmVersionUncached", "RowMajorTid"), + "SmVersionUncached", "RowMajorTid", "StoreDirectBlocked", + "StoreDirectStriped"), hasAncestor(namespaceDecl(hasName("cub"))))))) .bind("IntrinsicCall"), this); @@ -588,14 +589,14 @@ void CubDeviceLevelRule::removeRedundantTempVar(const CallExpr *CE) { void CubRule::registerMatcher(ast_matchers::MatchFinder &MF) { MF.addMatcher( typeLoc(loc(qualType(hasDeclaration(namedDecl(hasAnyName( - "WarpScan", "WarpReduce", "BlockScan", "BlockReduce", "BlockLoad")))))) + "WarpScan", "WarpReduce", "BlockScan", "BlockReduce")))))) .bind("TypeLoc"), this); MF.addMatcher( typedefDecl( hasType(hasCanonicalType(qualType(hasDeclaration(namedDecl(hasAnyName( - "WarpScan", "WarpReduce", "BlockScan", "BlockReduce", "BlockLoad"))))))) + "WarpScan", "WarpReduce", "BlockScan", "BlockReduce"))))))) .bind("TypeDefDecl"), this); @@ -729,8 +730,7 @@ void CubRule::processCubDeclStmt(const DeclStmt *DS) { ObjTypeStr.find("class cub::WarpReduce") == 0) { Repl = DpctGlobalInfo::getSubGroup(DRE); } else if (ObjTypeStr.find("class cub::BlockScan") == 0 || - ObjTypeStr.find("class cub::BlockReduce") == 0 || - ObjTypeStr.find("class cub::BlockLoad") == 0) { + ObjTypeStr.find("class cub::BlockReduce") == 0) { Repl = DpctGlobalInfo::getGroup(DRE); } else { continue; @@ -790,8 +790,7 @@ void CubRule::processCubTypeDef(const TypedefDecl *TD) { !(ObjTypeStr.find("class cub::WarpScan") == 0 || ObjTypeStr.find("class cub::WarpReduce") == 0 || ObjTypeStr.find("class cub::BlockScan") == 0 || - ObjTypeStr.find("class cub::BlockReduce") == 0 || - ObjTypeStr.find("class cub::BlockLoad") == 0)) { + ObjTypeStr.find("class cub::BlockReduce") == 0)) { DeleteFlag = false; break; } @@ -1497,8 +1496,7 @@ void CubRule::processCubMemberCall(const CXXMemberCallExpr *MC) { ObjTypeStr.find("class cub::WarpReduce") == 0) { processWarpLevelMemberCall(MC); } else if (ObjTypeStr.find("class cub::BlockScan") == 0 || - ObjTypeStr.find("class cub::BlockReduce") == 0 || - ObjTypeStr.find("class cub::BlockLoad") == 0) { + ObjTypeStr.find("class cub::BlockReduce") == 0) { processBlockLevelMemberCall(MC); } else { report(MC->getBeginLoc(), Diagnostics::API_NOT_MIGRATED, false, ObjTypeStr); @@ -1522,8 +1520,7 @@ void CubRule::processTypeLoc(const TypeLoc *TL) { MapNames::getClNamespace() + "sub_group", SM)); } else if (TypeName.find("class cub::BlockScan") == 0 || - TypeName.find("class cub::BlockReduce") == 0 || - TypeName.find("class cub::BlockLoad") == 0) { + TypeName.find("class cub::BlockReduce") == 0 ) { auto DeviceFuncDecl = DpctGlobalInfo::findAncestor(TL); if (DeviceFuncDecl && (DeviceFuncDecl->hasAttr() || DeviceFuncDecl->hasAttr())) { From b336ff2877c80253af69004755bff40e71ea21aa Mon Sep 17 00:00:00 2001 From: Abhilash Majumder <30946547+abhilash1910@users.noreply.github.com> Date: Thu, 11 Jul 2024 11:33:53 +0530 Subject: [PATCH 09/19] add store rewriter --- .../Rewriters/CUB/RewriterUtilityFunctions.cpp | 15 +++++++++++++++ 1 file changed, 15 insertions(+) diff --git a/clang/lib/DPCT/Rewriters/CUB/RewriterUtilityFunctions.cpp b/clang/lib/DPCT/Rewriters/CUB/RewriterUtilityFunctions.cpp index 3c31398a412d..824ee5cb5e8b 100644 --- a/clang/lib/DPCT/Rewriters/CUB/RewriterUtilityFunctions.cpp +++ b/clang/lib/DPCT/Rewriters/CUB/RewriterUtilityFunctions.cpp @@ -8,6 +8,7 @@ #include "CallExprRewriterCUB.h" #include "CallExprRewriterCommon.h" +#include "InclusionHeaders.h" using namespace clang::dpct; @@ -112,6 +113,20 @@ RewriterMap dpct::createUtilityFunctionsRewriterMap() { CALL(MapNames::getDpctNamespace() + "get_minor_version", makeDeviceStr()), LITERAL("10")))) + // cub::StoreDirectBlocked + HEADER_INSERT_FACTORY( + HeaderType::HT_DPCT_GROUP_Utils, + CALL_FACTORY_ENTRY( + "cub::StoreDirectBlocked", + CALL(MapNames::getDpctNamespace() + "group::store_blocked", NDITEM, + ARG(1), ARG(2)))) + // cub::StoreDirectStriped + HEADER_INSERT_FACTORY( + HeaderType::HT_DPCT_GROUP_Utils, + CALL_FACTORY_ENTRY( + "cub::StoreDirectStriped", + CALL(MapNames::getDpctNamespace() + "group::store_striped", NDITEM, + ARG(1), ARG(2)))) // cub::RowMajorTid MEMBER_CALL_FACTORY_ENTRY("cub::RowMajorTid", NDITEM, /*IsArrow=*/false, "get_local_linear_id")}; From 2550602002c7231698c946b39fc4249655941896 Mon Sep 17 00:00:00 2001 From: Abhilash Majumder <30946547+abhilash1910@users.noreply.github.com> Date: Thu, 11 Jul 2024 11:42:12 +0530 Subject: [PATCH 10/19] revert changes --- clang/lib/DPCT/ExprAnalysis.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/DPCT/ExprAnalysis.cpp b/clang/lib/DPCT/ExprAnalysis.cpp index 5c59742e474b..7c0a571f034d 100644 --- a/clang/lib/DPCT/ExprAnalysis.cpp +++ b/clang/lib/DPCT/ExprAnalysis.cpp @@ -1142,7 +1142,7 @@ void ExprAnalysis::analyzeType(TypeLoc TL, const Expr *CSCE, } } if (OS.str() != "cub::WarpScan" && OS.str() != "cub::WarpReduce" && - OS.str() != "cub::BlockReduce" && OS.str() != "cub::BlockScan" && OS.str() != "cub::BlockLoad") { + OS.str() != "cub::BlockReduce" && OS.str() != "cub::BlockScan" && OS.str()) { SR.setEnd(TSTL.getTemplateNameLoc()); } analyzeTemplateSpecializationType(TSTL); From 7efbc811bf108232a80951c0b3638eb143445f3d Mon Sep 17 00:00:00 2001 From: Abhilash Majumder <30946547+abhilash1910@users.noreply.github.com> Date: Thu, 11 Jul 2024 11:43:26 +0530 Subject: [PATCH 11/19] revert changes --- clang/lib/DPCT/ExprAnalysis.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/DPCT/ExprAnalysis.cpp b/clang/lib/DPCT/ExprAnalysis.cpp index 7c0a571f034d..347f931103e6 100644 --- a/clang/lib/DPCT/ExprAnalysis.cpp +++ b/clang/lib/DPCT/ExprAnalysis.cpp @@ -1142,7 +1142,7 @@ void ExprAnalysis::analyzeType(TypeLoc TL, const Expr *CSCE, } } if (OS.str() != "cub::WarpScan" && OS.str() != "cub::WarpReduce" && - OS.str() != "cub::BlockReduce" && OS.str() != "cub::BlockScan" && OS.str()) { + OS.str() != "cub::BlockReduce" && OS.str() != "cub::BlockScan") { SR.setEnd(TSTL.getTemplateNameLoc()); } analyzeTemplateSpecializationType(TSTL); From 5f7d10b6f28c09521de4d8ecbc7a392817fcacae Mon Sep 17 00:00:00 2001 From: Abhilash Majumder <30946547+abhilash1910@users.noreply.github.com> Date: Thu, 11 Jul 2024 11:51:04 +0530 Subject: [PATCH 12/19] fix issue --- clang/lib/DPCT/CUBAPIMigration.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/DPCT/CUBAPIMigration.cpp b/clang/lib/DPCT/CUBAPIMigration.cpp index 11ad50b17238..3cfcf02af8ee 100644 --- a/clang/lib/DPCT/CUBAPIMigration.cpp +++ b/clang/lib/DPCT/CUBAPIMigration.cpp @@ -613,7 +613,7 @@ void CubRule::registerMatcher(ast_matchers::MatchFinder &MF) { MF.addMatcher(cxxMemberCallExpr(has(memberExpr(member(hasAnyName( "InclusiveSum", "ExclusiveSum", "InclusiveScan", "ExclusiveScan", - "Reduce", "Sum", "Broadcast", "Scan", "Load"))))) + "Reduce", "Sum", "Broadcast", "Scan"))))) .bind("MemberCall"), this); From 3b0494f1fcb28870c92c8da7b7660f10cea3e1ff Mon Sep 17 00:00:00 2001 From: Abhilash Majumder <30946547+abhilash1910@users.noreply.github.com> Date: Thu, 11 Jul 2024 11:52:47 +0530 Subject: [PATCH 13/19] remove manual parsing --- clang/lib/DPCT/CUBAPIMigration.cpp | 11 ----------- 1 file changed, 11 deletions(-) diff --git a/clang/lib/DPCT/CUBAPIMigration.cpp b/clang/lib/DPCT/CUBAPIMigration.cpp index 3cfcf02af8ee..e564f6ce25c6 100644 --- a/clang/lib/DPCT/CUBAPIMigration.cpp +++ b/clang/lib/DPCT/CUBAPIMigration.cpp @@ -1319,17 +1319,6 @@ void CubRule::processBlockLevelMemberCall(const CXXMemberCallExpr *BlockMC) { CubParamAs << GroupOrWorkitem << In << OpRepl; Repl = NewFuncName + "(" + ParamList + ")"; emplaceTransformation(new ReplaceStmt(BlockMC, Repl)); - } else if (FuncName == "Load") { - - GroupOrWorkitem = DpctGlobalInfo::getItem(BlockMC); - NewFuncName = MapNames::getClNamespace() + "Load"; - const Expr *InData = FuncArgs[0]; - ExprAnalysis InEA(InData); - OpRepl = getOpRepl(nullptr); - CubParamAs << GroupOrWorkitem << InEA.getReplacedString() << OpRepl; - Repl = NewFuncName + "(" + ParamList + ")"; - emplaceTransformation(new ReplaceStmt(BlockMC, Repl)); - } } From d2ad0e5ec6660b18bd8ab57025a71466247f3ed9 Mon Sep 17 00:00:00 2001 From: Abhilash Majumder <30946547+abhilash1910@users.noreply.github.com> Date: Thu, 11 Jul 2024 13:47:30 +0530 Subject: [PATCH 14/19] Update APINames_CUB.inc --- clang/lib/DPCT/APINames_CUB.inc | 6 ++---- 1 file changed, 2 insertions(+), 4 deletions(-) diff --git a/clang/lib/DPCT/APINames_CUB.inc b/clang/lib/DPCT/APINames_CUB.inc index 35ff5b01857d..bd7c0dc6a4d5 100644 --- a/clang/lib/DPCT/APINames_CUB.inc +++ b/clang/lib/DPCT/APINames_CUB.inc @@ -122,8 +122,6 @@ ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Offset, Offset, fals ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Rotate, Rotate, false, NO_FLAG, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Up, Up, false, NO_FLAG, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Down, Down, false, NO_FLAG, P4, "Comment") -ENTRY_MEMBER_FUNCTION(cub::BlockLoad, cub::BlockLoad, Load, Load, true, NO_FLAG, P4, "Successful") - // Device Level ENTRY_MEMBER_FUNCTION(cub::DeviceAdjacentDifference, cub::DeviceAdjacentDifference, SubtractLeftCopy, SubtractLeftCopy, false, NO_FLAGE, P4, "Comment") @@ -192,10 +190,10 @@ ENTRY(cub::ThreadLoad, cub::ThreadLoad, true, NO_FLAG, P4, "Successful") ENTRY(cub::ThreadStore, cub::ThreadStore, true, NO_FLAG, P4, "Successful") ENTRY(cub::LoadDirectBlocked, cub::LoadDirectBlocked, false, NO_FLAG, P4, "Comment") ENTRY(cub::LoadDirectBlockedVectorized, cub::LoadDirectBlockedVectorized, false, NO_FLAG, P4, "Comment") -ENTRY(cub::StoreDirectBlocked, cub::StoreDirectBlocked, false, NO_FLAG, P4, "Comment") +ENTRY(cub::StoreDirectBlocked, cub::StoreDirectBlocked, true, NO_FLAG, P4, "Comment") ENTRY(cub::StoreDirectBlockedVectorized, cub::StoreDirectBlockedVectorized, false, NO_FLAG, P4, "Comment") ENTRY(cub::LoadDirectStriped, cub::LoadDirectStriped, false, NO_FLAG, P4, "Comment") -ENTRY(cub::StoreDirectStriped, cub::StoreDirectStriped, false, NO_FLAG, P4, "Comment") +ENTRY(cub::StoreDirectStriped, cub::StoreDirectStriped, true, NO_FLAG, P4, "Comment") ENTRY(cub::LoadDirectWarpStriped, cub::LoadDirectWarpStriped, false, NO_FLAG, P4, "Comment") ENTRY(cub::StoreDirectWarpStriped, cub::StoreDirectWarpStriped, false, NO_FLAG, P4, "Comment") From 9993f85fe56c043c12cff8569f000ae89d3a416b Mon Sep 17 00:00:00 2001 From: abhilash1910 Date: Thu, 11 Jul 2024 19:44:09 +0530 Subject: [PATCH 15/19] fix test --- clang/test/dpct/cub/blocklevel/blockload.cu | 83 --------------------- clang/test/dpct/cub/intrinsic/store.cu | 23 ++++++ 2 files changed, 23 insertions(+), 83 deletions(-) delete mode 100644 clang/test/dpct/cub/blocklevel/blockload.cu create mode 100644 clang/test/dpct/cub/intrinsic/store.cu diff --git a/clang/test/dpct/cub/blocklevel/blockload.cu b/clang/test/dpct/cub/blocklevel/blockload.cu deleted file mode 100644 index 927e806e6bf3..000000000000 --- a/clang/test/dpct/cub/blocklevel/blockload.cu +++ /dev/null @@ -1,83 +0,0 @@ -// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2 -// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2 -// RUN: dpct -in-root %S -out-root %T/blocklevel/blockload %S/blockload.cu --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only -// RUN: FileCheck --input-file %T/blocklevel/blockscan/blockload.dp.cpp --match-full-lines %s - -#include -#include - -#include -#include - -#define WARP_SIZE 32 - -const int N = 256; -const int BlockSize = 128; -const int ItemsPerThread = 4; - - -void init_data(int* data, int num) { - for(int i = 0; i < num; i++) - data[i] = i; -} -void verify_data(int* data, int num) { - return; -} -void print_data(int* data, int num) { - for (int i = 0; i < num; i++) { - std::cout << data[i] << " "; - } - std::cout << std::endl; -} - -//CHECK: void BlockLoadKernel(int* data, -//CHECK-NEXT: const sycl::nd_item<3> &item_ct1) { -//CHECK-EMPTY: -//CHECK-NEXT: int threadid = item_ct1.get_local_id(2); -//CHECK-EMPTY: -//CHECK-NEXT: int input = data[threadid]; -//CHECK-NEXT: int output = 0; -//CHECK-NEXT: output = sycl::load(item_ct1.get_group(), input, sycl::plus<>()); -//CHECK-NEXT: data[threadid] = output; -//CHECK-NEXT:} - -__global__ void BlockLoadKernel(int *d_data) -{ - // Specialize BlockLoad for a 1D block of 128 threads owning 4 integer items each - typedef cub::BlockLoad BlockLoad; - // Allocate shared memory for BlockLoad - __shared__ typename BlockLoad::TempStorage temp_storage; - // Load a segment of consecutive items that are blocked across threads - int thread_data[ItemsPerThread]; - int offset = threadIdx.x * ItemsPerThread; - BlockLoad(temp_storage).Load(d_data + offset, thread_data); - - // Print loaded data - printf("Thread %d loaded: %d %d %d %d\n", threadIdx.x, thread_data[0], thread_data[1], thread_data[2], thread_data[3]); -} - -int main() -{ - int h_data[N]; - init_data(h_data, N); - int *d_data; - cudaMalloc((void**)&d_data, N * sizeof(int)); - cudaMemcpy(d_data, h_data, N * sizeof(int), cudaMemcpyHostToDevice); - - //CHECK: q_ct1.parallel_for( - //CHECK-NEXT: sycl::nd_range<3>(GridSize * BlockSize, BlockSize), - //CHECK-NEXT: [=](sycl::nd_item<3> item_ct1) { - //CHECK-NEXT: BlockLoadKernel(dev_data, item_ct1); - //CHECK-NEXT: }); - - dim3 block(BlockSize); - dim3 grid((N + BlockSize - 1) / BlockSize); - - BlockLoadKernel<<>>(d_data); - cudaDeviceSynchronize(); - //verify_data(d_data, N); - - cudaFree(d_data); - - return 0; -} diff --git a/clang/test/dpct/cub/intrinsic/store.cu b/clang/test/dpct/cub/intrinsic/store.cu new file mode 100644 index 000000000000..0519babb0b82 --- /dev/null +++ b/clang/test/dpct/cub/intrinsic/store.cu @@ -0,0 +1,23 @@ +// UNSUPPORTED: cuda-8.0, cuda-9.0, cuda-9.1, cuda-9.2, cuda-10.0, cuda-10.1, cuda-10.2 +// UNSUPPORTED: v8.0, v9.0, v9.1, v9.2, v10.0, v10.1, v10.2 +// RUN: dpct -in-root %S -out-root %T/intrinsic/store %S/store.cu --cuda-include-path="%cuda-path/include" -- -std=c++14 -x cuda --cuda-host-only +// RUN: FileCheck --input-file %T/intrinsic/store/store.dp.cpp --match-full-lines %s +// RUN: %if build_lit %{icpx -c -fsycl %T/intrinsic/store/store.dp.cpp -o %T/intrinsic/store/store.dp.o %} + +// CHECK:#include +// CHECK:#include +// CHECK:#include +#include + +__global__ void TestStoreStriped(int *d_data) { + int thread_data[4]; + // CHECK: dpct::group::load_striped(item_ct1, d_data, thread_data); + cub::StoreDirectStriped<128>(threadIdx.x, d_data, thread_data); +} + + +__global__ void BlockedToStripedKernel(int *d_data) { + int thread_data[4]; + // CHECK: dpct::group::load_blocked(item_ct1, d_data, thread_data); + cub::StoreDirectBlocked(threadIdx.x, d_data, thread_data); +} \ No newline at end of file From a064b467727d27be6f9f5e043e4dbe43b441762a Mon Sep 17 00:00:00 2001 From: Abhilash Majumder <30946547+abhilash1910@users.noreply.github.com> Date: Thu, 8 Aug 2024 18:20:03 +0530 Subject: [PATCH 16/19] format fix --- clang/lib/DPCT/CUBAPIMigration.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/lib/DPCT/CUBAPIMigration.cpp b/clang/lib/DPCT/CUBAPIMigration.cpp index e564f6ce25c6..fe795bdb64f0 100644 --- a/clang/lib/DPCT/CUBAPIMigration.cpp +++ b/clang/lib/DPCT/CUBAPIMigration.cpp @@ -192,8 +192,8 @@ void CubIntrinsicRule::registerMatcher(ast_matchers::MatchFinder &MF) { "WarpId", "SyncStream", "CurrentDevice", "DeviceCount", "DeviceCountUncached", "DeviceCountCachedValue", "PtxVersion", "PtxVersionUncached", "SmVersion", - "SmVersionUncached", "RowMajorTid", "StoreDirectBlocked", - "StoreDirectStriped"), + "SmVersionUncached", "RowMajorTid", + "StoreDirectBlocked", "StoreDirectStriped"), hasAncestor(namespaceDecl(hasName("cub"))))))) .bind("IntrinsicCall"), this); @@ -1509,7 +1509,7 @@ void CubRule::processTypeLoc(const TypeLoc *TL) { MapNames::getClNamespace() + "sub_group", SM)); } else if (TypeName.find("class cub::BlockScan") == 0 || - TypeName.find("class cub::BlockReduce") == 0 ) { + TypeName.find("class cub::BlockReduce") == 0) { auto DeviceFuncDecl = DpctGlobalInfo::findAncestor(TL); if (DeviceFuncDecl && (DeviceFuncDecl->hasAttr() || DeviceFuncDecl->hasAttr())) { From 23735aa4066bdeb99701b4116b6899f982dd5025 Mon Sep 17 00:00:00 2001 From: Abhilash Majumder <30946547+abhilash1910@users.noreply.github.com> Date: Thu, 8 Aug 2024 18:20:46 +0530 Subject: [PATCH 17/19] add newline --- clang/test/dpct/cub/intrinsic/store.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/dpct/cub/intrinsic/store.cu b/clang/test/dpct/cub/intrinsic/store.cu index 0519babb0b82..0f89fdcc22fc 100644 --- a/clang/test/dpct/cub/intrinsic/store.cu +++ b/clang/test/dpct/cub/intrinsic/store.cu @@ -20,4 +20,4 @@ __global__ void BlockedToStripedKernel(int *d_data) { int thread_data[4]; // CHECK: dpct::group::load_blocked(item_ct1, d_data, thread_data); cub::StoreDirectBlocked(threadIdx.x, d_data, thread_data); -} \ No newline at end of file +} From 5fc94c64390039df1b540231db755642298985d2 Mon Sep 17 00:00:00 2001 From: Abhilash Majumder <30946547+abhilash1910@users.noreply.github.com> Date: Tue, 13 Aug 2024 18:46:19 +0530 Subject: [PATCH 18/19] fix include file --- clang/lib/DPCT/APINames_CUB.inc | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/clang/lib/DPCT/APINames_CUB.inc b/clang/lib/DPCT/APINames_CUB.inc index bd7c0dc6a4d5..880ed9f92a1a 100644 --- a/clang/lib/DPCT/APINames_CUB.inc +++ b/clang/lib/DPCT/APINames_CUB.inc @@ -107,6 +107,7 @@ ENTRY_MEMBER_FUNCTION(cub::BlockExchange, cub::BlockExchange, ScatterToStripedFl ENTRY_MEMBER_FUNCTION(cub::BlockHistogram, cub::BlockHistogram, InitHistogram, InitHistogram, false, NO_FLAG, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::BlockHistogram, cub::BlockHistogram, Histogram, Histogram, false, NO_FLAG, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::BlockHistogram, cub::BlockHistogram, Composite, Composite, false, NO_FLAG, P4, "Comment") +ENTRY_MEMBER_FUNCTION(cub::BlockLoad, cub::BlockLoad, Load, Load, false, NO_FLAG, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::BlockStore, cub::BlockStore, Store, Store, false, NO_FLAG, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::BlockRadixSort, cub::BlockRadixSort, Sort, Sort, false, NO_FLAG, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::BlockRadixSort, cub::BlockRadixSort, SortDescending, SortDescending, false, NO_FLAG, P4, "Comment") @@ -190,10 +191,10 @@ ENTRY(cub::ThreadLoad, cub::ThreadLoad, true, NO_FLAG, P4, "Successful") ENTRY(cub::ThreadStore, cub::ThreadStore, true, NO_FLAG, P4, "Successful") ENTRY(cub::LoadDirectBlocked, cub::LoadDirectBlocked, false, NO_FLAG, P4, "Comment") ENTRY(cub::LoadDirectBlockedVectorized, cub::LoadDirectBlockedVectorized, false, NO_FLAG, P4, "Comment") -ENTRY(cub::StoreDirectBlocked, cub::StoreDirectBlocked, true, NO_FLAG, P4, "Comment") +ENTRY(cub::StoreDirectBlocked, cub::StoreDirectBlocked, true, NO_FLAG, P4, "Successful") ENTRY(cub::StoreDirectBlockedVectorized, cub::StoreDirectBlockedVectorized, false, NO_FLAG, P4, "Comment") ENTRY(cub::LoadDirectStriped, cub::LoadDirectStriped, false, NO_FLAG, P4, "Comment") -ENTRY(cub::StoreDirectStriped, cub::StoreDirectStriped, true, NO_FLAG, P4, "Comment") +ENTRY(cub::StoreDirectStriped, cub::StoreDirectStriped, true, NO_FLAG, P4, "Successful") ENTRY(cub::LoadDirectWarpStriped, cub::LoadDirectWarpStriped, false, NO_FLAG, P4, "Comment") ENTRY(cub::StoreDirectWarpStriped, cub::StoreDirectWarpStriped, false, NO_FLAG, P4, "Comment") From fd9c641828f460cdadc6e6ceaa50ad961a2891fe Mon Sep 17 00:00:00 2001 From: Abhilash Majumder <30946547+abhilash1910@users.noreply.github.com> Date: Tue, 13 Aug 2024 19:44:14 +0530 Subject: [PATCH 19/19] fix blank line --- clang/lib/DPCT/APINames_CUB.inc | 1 + 1 file changed, 1 insertion(+) diff --git a/clang/lib/DPCT/APINames_CUB.inc b/clang/lib/DPCT/APINames_CUB.inc index 880ed9f92a1a..30e3c58d8ceb 100644 --- a/clang/lib/DPCT/APINames_CUB.inc +++ b/clang/lib/DPCT/APINames_CUB.inc @@ -124,6 +124,7 @@ ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Rotate, Rotate, fals ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Up, Up, false, NO_FLAG, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::BlockShuffle, cub::BlockShuffle, Down, Down, false, NO_FLAG, P4, "Comment") + // Device Level ENTRY_MEMBER_FUNCTION(cub::DeviceAdjacentDifference, cub::DeviceAdjacentDifference, SubtractLeftCopy, SubtractLeftCopy, false, NO_FLAGE, P4, "Comment") ENTRY_MEMBER_FUNCTION(cub::DeviceAdjacentDifference, cub::DeviceAdjacentDifference, SubtractLeft, SubtractLeft, false, NO_FLAGE, P4, "Comment")