From 2bf3a98d8323b0b28e4eb5c5d96e799b024080e2 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Mon, 23 Oct 2017 16:47:40 -0400 Subject: [PATCH 1/5] replace __hcc_workweek__ with HC_FEATURE_PRINTF flag [ROCm/hip commit: 5b9ce032d640e12928d66ea9aa7efb804efe3f1c] --- projects/hip/src/hip_hcc.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/src/hip_hcc.cpp b/projects/hip/src/hip_hcc.cpp index 2d67c31fe7..ae53d7ae45 100644 --- a/projects/hip/src/hip_hcc.cpp +++ b/projects/hip/src/hip_hcc.cpp @@ -616,7 +616,7 @@ void ihipDevice_t::locked_reset() //FIXME - Calling am_memtracker_reset is really bad since it destroyed all buffers allocated by the HCC runtime as well //such as the printf buffer. Re-initialze the printf buffer as a workaround for now. -#if (__hcc_workweek__ >= 17423) +#ifdef HC_FEATURE_PRINTF Kalmar::getContext()->initPrintfBuffer(); #endif }; From ccc193cf9877eae8d803b57e8c036234082f7326 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Mon, 23 Oct 2017 16:59:14 -0400 Subject: [PATCH 2/5] add a printf device function [ROCm/hip commit: f8bdb9f688c1633e9cf1f2290173a8862b33fdbc] --- .../hip/include/hip/hcc_detail/hip_runtime.h | 20 +++++++++++++++++++ 1 file changed, 20 insertions(+) diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h index 379fc05f5b..07c7a0d4b3 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime.h @@ -50,10 +50,16 @@ THE SOFTWARE. #include +// define HIP_ENABLE_PRINTF to enable printf +#ifdef HIP_ENABLE_PRINTF + #define HCC_ENABLE_ACCELERATOR_PRINTF 1 +#endif + //--- // Remainder of this file only compiles with HCC #if defined __HCC__ #include +#include "hc_printf.hpp" //TODO-HCC-GL - change this to typedef. //typedef grid_launch_parm hipLaunchParm ; @@ -420,6 +426,20 @@ static inline __device__ void* memset(void* ptr, int val, size_t size) } +#ifdef __HCC_ACCELERATOR__ + +#ifdef HC_FEATURE_PRINTF +template +static inline __device__ void printf(const char* format, All... all) { + hc::printf(format, all...); +} +#else +template +static inline __device__ void printf(const char* format, All... all) { } +#endif + +#endif + #define __syncthreads() hc_barrier(CLK_LOCAL_MEM_FENCE) From 76f70bd60e53eb042fc9780dfee8da42df65f182 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Mon, 23 Oct 2017 18:31:13 -0400 Subject: [PATCH 3/5] add an unit test for printf [ROCm/hip commit: b7b00fb281b5b81e5b6732d269f3620f6a040e3b] --- .../hip/tests/src/kernel/hipPrintfKernel.cpp | 38 +++++++++++++++++++ 1 file changed, 38 insertions(+) create mode 100644 projects/hip/tests/src/kernel/hipPrintfKernel.cpp diff --git a/projects/hip/tests/src/kernel/hipPrintfKernel.cpp b/projects/hip/tests/src/kernel/hipPrintfKernel.cpp new file mode 100644 index 0000000000..482098fd54 --- /dev/null +++ b/projects/hip/tests/src/kernel/hipPrintfKernel.cpp @@ -0,0 +1,38 @@ +/* +Copyright (c) 2015-2017 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. +*/ + +/* HIT_START + * BUILD: %t %s ../test_common.cpp + * RUN: %t + * HIT_END + */ + +#define HIP_ENABLE_PRINTF + +#include"test_common.h" + +__global__ void run_printf(hipLaunchParm lp){ + printf("Hello World\n"); +} + +int main(){ +hipLaunchKernel(HIP_KERNEL_NAME(run_printf), dim3(1), dim3(1), 0, 0); +hipDeviceSynchronize(); +passed(); +} From 638f62b0c49839cb19563c47f942a9c614abd70a Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Tue, 24 Oct 2017 00:55:41 -0400 Subject: [PATCH 4/5] printf support for module API [ROCm/hip commit: 1ddee10c2f9a6e8a5f9634c78c9c5c8b005093c4] --- projects/hip/src/hip_module.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index 74b0092b84..85c88a1826 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -259,6 +259,11 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable); CHECKLOG_HSA(status, hipErrorNotInitialized); + status = hsa_executable_agent_global_variable_define((*module)->executable, agent, + "_ZN2hc13printf_bufferE", + Kalmar::getContext()->getPrintfBufferPointerVA()); + CHECKLOG_HSA(status, hipErrorNotInitialized); + status = hsa_executable_load_code_object((*module)->executable, agent, (*module)->object, NULL); CHECKLOG_HSA(status, hipErrorNotInitialized); From c595b8f660b42409032b3bdbf3273b3f2af5ed54 Mon Sep 17 00:00:00 2001 From: Siu Chi Chan Date: Wed, 25 Oct 2017 12:00:02 -0400 Subject: [PATCH 5/5] add HC_FEATURE_PRINTF around the printf buffer definition [ROCm/hip commit: d91a4f5bd60cf5baaf4b1fabc369c4a8d37b7a91] --- projects/hip/src/hip_module.cpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index 85c88a1826..962b6084f1 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -259,10 +259,12 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable); CHECKLOG_HSA(status, hipErrorNotInitialized); +#ifdef HC_FEATURE_PRINTF status = hsa_executable_agent_global_variable_define((*module)->executable, agent, "_ZN2hc13printf_bufferE", Kalmar::getContext()->getPrintfBufferPointerVA()); CHECKLOG_HSA(status, hipErrorNotInitialized); +#endif status = hsa_executable_load_code_object((*module)->executable, agent, (*module)->object, NULL); CHECKLOG_HSA(status, hipErrorNotInitialized);