From 643a8bcf5bff879a706c741a8492aa83060a1271 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 16 Oct 2019 13:08:11 +0300 Subject: [PATCH] [HIPIFY][CUB][#1460] Implement cubFunctionTemplateDecl matcher + Add cub_02.cu test + Partial fixes #1460 [ROCm/clr commit: 5555d46e6663b303cad5f8ab6178bc2bf5522fb8] --- .../hipamd/hipify-clang/src/HipifyAction.cpp | 44 ++++++++++++ .../hipamd/hipify-clang/src/HipifyAction.h | 1 + .../unit_tests/libraries/CUB/cub_02.cu | 70 +++++++++++++++++++ 3 files changed, 115 insertions(+) create mode 100644 projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/CUB/cub_02.cu diff --git a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp index 930f3ec8c5..60288e7500 100644 --- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp +++ b/projects/clr/hipamd/hipify-clang/src/HipifyAction.cpp @@ -61,6 +61,7 @@ const StringRef sCudaLaunchKernel = "cudaLaunchKernel"; const StringRef sCudaHostFuncCall = "cudaHostFuncCall"; const StringRef sCudaDeviceFuncCall = "cudaDeviceFuncCall"; const StringRef sCubNamespacePrefix = "cubNamespacePrefix"; +const StringRef sCubFunctionTemplateDecl = "cubFunctionTemplateDecl"; std::set DeviceSymbolFunctions0 { {sCudaMemcpyToSymbol}, @@ -449,6 +450,41 @@ bool HipifyAction::cubNamespacePrefix(const mat::MatchFinder::MatchResult &Resul return false; } +bool HipifyAction::cubFunctionTemplateDecl(const mat::MatchFinder::MatchResult &Result) { + if (auto *decl = Result.Nodes.getNodeAs(sCubFunctionTemplateDecl)) { + auto *Tparams = decl->getTemplateParameters(); + bool ret = false; + for (size_t I = 0; I < Tparams->size(); ++I) { + const clang::ValueDecl *valueDecl = dyn_cast(Tparams->getParam(I)); + if (!valueDecl) continue; + clang::QualType QT = valueDecl->getType(); + auto *t = QT.getTypePtr(); + if (!t) continue; + const clang::ElaboratedType *et = t->getAs(); + if (!et) continue; + const clang::NestedNameSpecifier *nns = et->getQualifier(); + if (!nns) continue; + const clang::NamespaceDecl *nsd = nns->getAsNamespace(); + if (!nsd) continue; + const clang::SourceRange sr = valueDecl->getSourceRange(); + clang::SourceLocation sl(sr.getBegin()); + clang::SourceLocation end(sr.getEnd()); + auto &SM = getCompilerInstance().getSourceManager(); + size_t length = SM.getCharacterData(end) - SM.getCharacterData(sl); + StringRef sfull = StringRef(SM.getCharacterData(sl), length); + std::string name = nsd->getDeclName().getAsString(); + size_t offset = sfull.find(name); + if (offset > 0) { + sl = sl.getLocWithOffset(offset); + } + FindAndReplace(name, sl, CUDA_CUB_TYPE_NAME_MAP); + ret = true; + } + return ret; + } + return false; +} + bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result) { if (auto *call = Result.Nodes.getNodeAs(sCudaHostFuncCall)) { if (!call->getNumArgs()) return false; @@ -555,6 +591,13 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi ).bind(sCubNamespacePrefix), this ); + // TODO: Maybe worth to make it more concrete based on final cubFunctionTemplateDecl + Finder->addMatcher( + mat::functionTemplateDecl( + mat::isExpansionInMainFile() + ).bind(sCubFunctionTemplateDecl), + this + ); // Ownership is transferred to the caller. return Finder->newASTConsumer(); } @@ -668,4 +711,5 @@ void HipifyAction::run(const mat::MatchFinder::MatchResult &Result) { if (cudaHostFuncCall(Result)) return; if (cudaDeviceFuncCall(Result)) return; if (cubNamespacePrefix(Result)) return; + if (cubFunctionTemplateDecl(Result)) return; } diff --git a/projects/clr/hipamd/hipify-clang/src/HipifyAction.h b/projects/clr/hipamd/hipify-clang/src/HipifyAction.h index 78532449ee..38a0deba9e 100644 --- a/projects/clr/hipamd/hipify-clang/src/HipifyAction.h +++ b/projects/clr/hipamd/hipify-clang/src/HipifyAction.h @@ -73,6 +73,7 @@ public: bool cudaDeviceFuncCall(const mat::MatchFinder::MatchResult &Result); bool cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result); bool cubNamespacePrefix(const mat::MatchFinder::MatchResult &Result); + bool cubFunctionTemplateDecl(const mat::MatchFinder::MatchResult &Result); // Called by the preprocessor for each include directive during the non-raw lexing pass. void InclusionDirective(clang::SourceLocation hash_loc, const clang::Token &include_token, diff --git a/projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/CUB/cub_02.cu b/projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/CUB/cub_02.cu new file mode 100644 index 0000000000..aff5def3fa --- /dev/null +++ b/projects/clr/hipamd/tests/hipify-clang/unit_tests/libraries/CUB/cub_02.cu @@ -0,0 +1,70 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +// CHECK: #include +#include +// CHECK: #include +#include +// CHECK: #include +#include + +#include + +template +__global__ void sort(const T* data_in, T* data_out){ + // CHECK: typedef hipcub::BlockLoad BlockLoadT; + typedef cub::BlockLoad BlockLoadT; + // CHECK: typedef hipcub::BlockRadixSort BlockRadixSortT; + typedef cub::BlockRadixSort BlockRadixSortT; + // CHECK: typedef hipcub::BlockStore BlockStoreT; + typedef cub::BlockStore BlockStoreT; + __shared__ union { + typename BlockLoadT::TempStorage load; + typename BlockRadixSortT::TempStorage sort; + typename BlockStoreT::TempStorage store; + } tmp_storage; + T items[ITEMS_PER_THREAD]; + BlockLoadT(tmp_storage.load).Load(data_in + blockIdx.x * BLOCK_WIDTH * ITEMS_PER_THREAD, items); + __syncthreads(); + BlockRadixSortT(tmp_storage.sort).Sort(items); + __syncthreads(); + BlockStoreT(tmp_storage.store).Store(data_out + blockIdx.x * BLOCK_WIDTH * ITEMS_PER_THREAD, items); +} + +int main() { + double* d_gpu = NULL; + double* result_gpu = NULL; + double* data_sorted = new double[1000*4096]; + // Allocate memory on the GPU + // CHECK: hipMalloc(&d_gpu, 1000*4096 * sizeof(double)); + cudaMalloc(&d_gpu, 1000*4096 * sizeof(double)); + // CHECK: hipMalloc(&result_gpu, 1000*4096 * sizeof(double)); + cudaMalloc(&result_gpu, 1000*4096 * sizeof(double)); + // CHECK: hiprandGenerator_t gen; + curandGenerator_t gen; + // Create generator + // CHECK: hiprandCreateGenerator(&gen, HIPRAND_RNG_PSEUDO_DEFAULT); + curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_DEFAULT); + // Fill array with random numbers + // CHECK: hiprandGenerateNormalDouble(gen, d_gpu, 1000*4096, 0.0, 1.0); + curandGenerateNormalDouble(gen, d_gpu, 1000*4096, 0.0, 1.0); + // Destroy generator + // CHECK: hiprandDestroyGenerator(gen); + curandDestroyGenerator(gen); + // Sort data + // TODO: Substitution of cub namespace in CUDAKernelCallExpr + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(sort<512, 8, cub::BLOCK_LOAD_TRANSPOSE, cub::BLOCK_STORE_TRANSPOSE>), dim3(1000), dim3(512), 0, 0, d_gpu, result_gpu); + sort<512, 8, cub::BLOCK_LOAD_TRANSPOSE, cub::BLOCK_STORE_TRANSPOSE><<<1000, 512>>>(d_gpu, result_gpu); + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(sort<256, 16, cub::BLOCK_LOAD_DIRECT, cub::BLOCK_STORE_DIRECT>), dim3(1000), dim3(256), 0, 0, d_gpu, result_gpu); + sort<256, 16, cub::BLOCK_LOAD_DIRECT, cub::BLOCK_STORE_DIRECT><<<1000, 256>>>(d_gpu, result_gpu); + // CHECK: hipMemcpy(data_sorted, result_gpu, 1000*4096*sizeof(double), hipMemcpyDeviceToHost); + cudaMemcpy(data_sorted, result_gpu, 1000*4096*sizeof(double), cudaMemcpyDeviceToHost); + // Write the sorted data to standard out + for (int i = 0; i < 4095; ++i) { + std::cout << data_sorted[i] << ", "; + } + std::cout << data_sorted[4095] << std::endl; +}