From c7a2b4e492785e38e5ae7ba5df282728275feaeb Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 23 May 2019 12:47:08 +0300 Subject: [PATCH] [HIPIFY] Add device texture fetch functions support + Add a corresponding reverse engineered sample tex2dKernel with texture template --- hipify-clang/src/CUDA2HIP.cpp | 1 + hipify-clang/src/HipifyAction.cpp | 3 ++ .../11_texture_driver/tex2dKernel.cpp | 36 +++++++++++++++++++ 3 files changed, 40 insertions(+) create mode 100644 tests/hipify-clang/unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp diff --git a/hipify-clang/src/CUDA2HIP.cpp b/hipify-clang/src/CUDA2HIP.cpp index 9a26a051ca..1e530745e6 100644 --- a/hipify-clang/src/CUDA2HIP.cpp +++ b/hipify-clang/src/CUDA2HIP.cpp @@ -33,6 +33,7 @@ const std::map CUDA_INCLUDE_MAP{ {"driver_types.h", {"hip/driver_types.h", "", CONV_INCLUDE, API_RUNTIME}}, {"cuda_fp16.h", {"hip/hip_fp16.h", "", CONV_INCLUDE, API_RUNTIME}}, {"cuda_texture_types.h", {"hip/hip_texture_types.h", "", CONV_INCLUDE, API_RUNTIME}}, + {"texture_fetch_functions.h", {"", "", CONV_INCLUDE, API_RUNTIME}}, {"vector_types.h", {"hip/hip_vector_types.h", "", CONV_INCLUDE, API_RUNTIME}}, {"cuda_profiler_api.h", {"hip/hip_profile.h", "", CONV_INCLUDE, API_RUNTIME}}, // cuComplex includes diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp index 5ea83c8376..241ca7ecae 100644 --- a/hipify-clang/src/HipifyAction.cpp +++ b/hipify-clang/src/HipifyAction.cpp @@ -189,6 +189,9 @@ bool HipifyAction::Exclude(const hipCounter & hipToken) { } return false; case CONV_INCLUDE: + if (hipToken.hipName.empty()) { + return true; + } switch (hipToken.apiType) { case API_RAND: if (hipToken.hipName == "hiprand_kernel.h") { diff --git a/tests/hipify-clang/unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp b/tests/hipify-clang/unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp new file mode 100644 index 0000000000..d5dffd0b09 --- /dev/null +++ b/tests/hipify-clang/unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp @@ -0,0 +1,36 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +/* +Copyright (c) 2015-present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +// CHECK: #include +#include +// CHECK-NOT: #include +#include + +// CHECK: extern texture tex; +extern texture tex; + +extern "C" __global__ void tex2dKernel(float* outputData, int width, int height) { + int x = blockDim.x * blockIdx.x + threadIdx.x; + int y = blockDim.y * blockIdx.y + threadIdx.y; + outputData[y * width + x] = tex2D(tex, x, y); +}