From 5a48bc2168dcf3713acb9b3a2abf3b9ff35a93da Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 1 Oct 2018 15:28:34 +0530 Subject: [PATCH 1/9] Fixed image width for linear resource type texture [ROCm/hip commit: 0b8fc1d66a26054f0c7d9d4e4b559a60b31cb96e] --- projects/hip/src/hip_texture.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/projects/hip/src/hip_texture.cpp b/projects/hip/src/hip_texture.cpp index 82bc5b5d8a..85fb006d43 100644 --- a/projects/hip/src/hip_texture.cpp +++ b/projects/hip/src/hip_texture.cpp @@ -198,6 +198,7 @@ bool getHipTextureObject(hipTextureObject_t* pTexObject, hsa_ext_image_t& image, return true; } +size_t channelFormatKindSize[] = {sizeof(unsigned),sizeof(signed),sizeof(float),sizeof(unsigned)}; // Texture Object APIs hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResourceDesc* pResDesc, const hipTextureDesc* pTexDesc, @@ -263,7 +264,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou break; case hipResourceTypeLinear: devPtr = pResDesc->res.linear.devPtr; - imageDescriptor.width = pResDesc->res.linear.sizeInBytes; + imageDescriptor.width = pResDesc->res.linear.sizeInBytes/channelFormatKindSize[pResDesc->res.linear.desc.f]; imageDescriptor.height = 1; imageDescriptor.depth = 0; imageDescriptor.array_size = 0; From f201230ad36fdec7a99a9e0615cf6be3600fc8b5 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Fri, 28 Sep 2018 11:05:27 -0400 Subject: [PATCH 2/9] Add HIPCC_COMPILE_FLAGS_APPEND [ROCm/hip commit: dcc85be374668398502e2e6e8bd6d4887798852b] --- projects/hip/bin/hipcc | 15 ++++++++++++--- 1 file changed, 12 insertions(+), 3 deletions(-) diff --git a/projects/hip/bin/hipcc b/projects/hip/bin/hipcc index 620a53671b..9821a6ceb2 100755 --- a/projects/hip/bin/hipcc +++ b/projects/hip/bin/hipcc @@ -53,6 +53,9 @@ sub parse_config_file { $verbose = $ENV{'HIPCC_VERBOSE'} // 0; # Verbose: 0x1=commands, 0x2=paths, 0x4=hipcc args +$HIPCC_COMPILE_FLAGS_APPEND=$ENV{'HIPCC_COMPILE_FLAGS_APPEND'}; +$HIPCC_LINK_FLAGS_APPEND=$ENV{'HIPCC_LINK_FLAGS_APPEND'}; + $HIP_PATH=$ENV{'HIP_PATH'} // dirname (dirname $0); # use parent directory of hipcc $HIP_VDI_HOME=$ENV{'HIP_VDI_HOME'}; $HIP_CLANG_PATH=$ENV{'HIP_CLANG_PATH'}; @@ -723,7 +726,16 @@ if ($needHipHcc) { # Reason is that NVCC uses the file extension to determine whether to compile in CUDA mode or # pass-through CPP mode. +if ($HIP_PLATFORM eq "clang") { + $HIPLDFLAGS .= " -lgcc_s -lgcc -lpthread -lm"; +} +if ($HIPCC_COMPILE_FLAGS_APPEND) { + $HIPCXXFLAGS .= " $HIPCC_COMPILE_FLAGS_APPEND"; +} +if ($HIPCC_LINK_FLAGS_APPEND) { + $HIPLDFLAGS .= " $HIPCC_LINK_FLAGS_APPEND"; +} my $CMD="$HIPCC"; if ($needCXXFLAGS) { @@ -733,9 +745,6 @@ if ($needLDFLAGS and not $compileOnly) { $CMD .= " $HIPLDFLAGS"; } $CMD .= " $toolArgs"; -if ($needLDFLAGS and not $compileOnly and $HIP_PLATFORM eq "clang") { - $CMD .= " -lgcc_s -lgcc -lpthread -lm"; -} if ($verbose & 0x1) { print "hipcc-cmd: ", $CMD, "\n"; From 8401146cd6db5d72e34300e61750a505c4836d4e Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 3 Oct 2018 12:07:38 +0530 Subject: [PATCH 3/9] Corrected the width calculation logic to accomodate multi channels [ROCm/hip commit: 5a3d89e99c3006f08a7eb631e9da6912daf52330] --- projects/hip/src/hip_texture.cpp | 3 +-- 1 file changed, 1 insertion(+), 2 deletions(-) diff --git a/projects/hip/src/hip_texture.cpp b/projects/hip/src/hip_texture.cpp index 85fb006d43..521e0e24de 100644 --- a/projects/hip/src/hip_texture.cpp +++ b/projects/hip/src/hip_texture.cpp @@ -198,7 +198,6 @@ bool getHipTextureObject(hipTextureObject_t* pTexObject, hsa_ext_image_t& image, return true; } -size_t channelFormatKindSize[] = {sizeof(unsigned),sizeof(signed),sizeof(float),sizeof(unsigned)}; // Texture Object APIs hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResourceDesc* pResDesc, const hipTextureDesc* pTexDesc, @@ -264,7 +263,7 @@ hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject, const hipResou break; case hipResourceTypeLinear: devPtr = pResDesc->res.linear.devPtr; - imageDescriptor.width = pResDesc->res.linear.sizeInBytes/channelFormatKindSize[pResDesc->res.linear.desc.f]; + imageDescriptor.width = pResDesc->res.linear.sizeInBytes/((pResDesc->res.linear.desc.x + pResDesc->res.linear.desc.y + pResDesc->res.linear.desc.z + pResDesc->res.linear.desc.w)/8); imageDescriptor.height = 1; imageDescriptor.depth = 0; imageDescriptor.array_size = 0; From 188df279a9343bb7cb7a388d0f69c95605093394 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Thu, 4 Oct 2018 21:47:28 +0000 Subject: [PATCH 4/9] Replace IRIF fences with atomic_work_item_fence [ROCm/hip commit: 0787f74ac28527bd2e5c490bb8b6e31eb84861ed] --- .../include/hip/hcc_detail/device_functions.h | 74 ------------------- .../hip/hcc_detail/device_library_decls.h | 40 +++++----- 2 files changed, 22 insertions(+), 92 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/device_functions.h b/projects/hip/include/hip/hcc_detail/device_functions.h index b319f26e03..42927e3246 100644 --- a/projects/hip/include/hip/hcc_detail/device_functions.h +++ b/projects/hip/include/hip/hcc_detail/device_functions.h @@ -752,80 +752,6 @@ void *__amdgcn_get_dynamicgroupbaseptr() { return __get_dynamicgroupbaseptr(); } - - -// hip.amdgcn.bc - sync threads -#define __CLK_LOCAL_MEM_FENCE 0x01 -typedef unsigned __cl_mem_fence_flags; - -typedef enum __memory_scope { - __memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM, - __memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP, - __memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE, - __memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES, - __memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP -} __memory_scope; - -// enum values aligned with what clang uses in EmitAtomicExpr() -typedef enum __memory_order -{ - __memory_order_relaxed = __ATOMIC_RELAXED, - __memory_order_acquire = __ATOMIC_ACQUIRE, - __memory_order_release = __ATOMIC_RELEASE, - __memory_order_acq_rel = __ATOMIC_ACQ_REL, - __memory_order_seq_cst = __ATOMIC_SEQ_CST -} __memory_order; - -__device__ -inline -static void -__atomic_work_item_fence(__cl_mem_fence_flags flags, __memory_order order, __memory_scope scope) -{ - // We're tying global-happens-before and local-happens-before together as does HSA - if (order != __memory_order_relaxed) { - switch (scope) { - case __memory_scope_work_item: - break; - case __memory_scope_sub_group: - switch (order) { - case __memory_order_relaxed: break; - case __memory_order_acquire: __llvm_fence_acq_sg(); break; - case __memory_order_release: __llvm_fence_rel_sg(); break; - case __memory_order_acq_rel: __llvm_fence_ar_sg(); break; - case __memory_order_seq_cst: __llvm_fence_sc_sg(); break; - } - break; - case __memory_scope_work_group: - switch (order) { - case __memory_order_relaxed: break; - case __memory_order_acquire: __llvm_fence_acq_wg(); break; - case __memory_order_release: __llvm_fence_rel_wg(); break; - case __memory_order_acq_rel: __llvm_fence_ar_wg(); break; - case __memory_order_seq_cst: __llvm_fence_sc_wg(); break; - } - break; - case __memory_scope_device: - switch (order) { - case __memory_order_relaxed: break; - case __memory_order_acquire: __llvm_fence_acq_dev(); break; - case __memory_order_release: __llvm_fence_rel_dev(); break; - case __memory_order_acq_rel: __llvm_fence_ar_dev(); break; - case __memory_order_seq_cst: __llvm_fence_sc_dev(); break; - } - break; - case __memory_scope_all_svm_devices: - switch (order) { - case __memory_order_relaxed: break; - case __memory_order_acquire: __llvm_fence_acq_sys(); break; - case __memory_order_release: __llvm_fence_rel_sys(); break; - case __memory_order_acq_rel: __llvm_fence_ar_sys(); break; - case __memory_order_seq_cst: __llvm_fence_sc_sys(); break; - } - break; - } - } -} - // Memory Fence Functions __device__ inline diff --git a/projects/hip/include/hip/hcc_detail/device_library_decls.h b/projects/hip/include/hip/hcc_detail/device_library_decls.h index d64877d50a..2a0b968fe3 100644 --- a/projects/hip/include/hip/hcc_detail/device_library_decls.h +++ b/projects/hip/include/hip/hcc_detail/device_library_decls.h @@ -65,26 +65,30 @@ extern "C" __device__ __attribute__((const)) float __ocml_fmax_f32(float, float) __device__ inline static __local void* __to_local(unsigned x) { return (__local void*)x; } #endif //__HIP_DEVICE_COMPILE__ -// __llvm_fence* functions from device-libs/irif/src/fence.ll -extern "C" __device__ void __llvm_fence_acq_sg(void); -extern "C" __device__ void __llvm_fence_acq_wg(void); -extern "C" __device__ void __llvm_fence_acq_dev(void); -extern "C" __device__ void __llvm_fence_acq_sys(void); +// Using hip.amdgcn.bc - sync threads +#define __CLK_LOCAL_MEM_FENCE 0x01 +typedef unsigned __cl_mem_fence_flags; -extern "C" __device__ void __llvm_fence_rel_sg(void); -extern "C" __device__ void __llvm_fence_rel_wg(void); -extern "C" __device__ void __llvm_fence_rel_dev(void); -extern "C" __device__ void __llvm_fence_rel_sys(void); +typedef enum __memory_scope { + __memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM, + __memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP, + __memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE, + __memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES, + __memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP +} __memory_scope; -extern "C" __device__ void __llvm_fence_ar_sg(void); -extern "C" __device__ void __llvm_fence_ar_wg(void); -extern "C" __device__ void __llvm_fence_ar_dev(void); -extern "C" __device__ void __llvm_fence_ar_sys(void); +// enum values aligned with what clang uses in EmitAtomicExpr() +typedef enum __memory_order +{ + __memory_order_relaxed = __ATOMIC_RELAXED, + __memory_order_acquire = __ATOMIC_ACQUIRE, + __memory_order_release = __ATOMIC_RELEASE, + __memory_order_acq_rel = __ATOMIC_ACQ_REL, + __memory_order_seq_cst = __ATOMIC_SEQ_CST +} __memory_order; - -extern "C" __device__ void __llvm_fence_sc_sg(void); -extern "C" __device__ void __llvm_fence_sc_wg(void); -extern "C" __device__ void __llvm_fence_sc_dev(void); -extern "C" __device__ void __llvm_fence_sc_sys(void); +// Linked from hip.amdgcn.bc +extern "C" __device__ void +__atomic_work_item_fence(__cl_mem_fence_flags, __memory_order, __memory_scope); #endif From b35590db1890af2f535a5db251086e7cff35c503 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 5 Oct 2018 16:23:08 +0530 Subject: [PATCH 5/9] Implementation for stream priority - Requires ROCm 1.9.x or higher - Requires HCC with PR#886 merged Change-Id: Id7c95ea091ee610e80c9ad815f1cb989cba570ca [ROCm/hip commit: da64156fb2d7ba42823b36e91fa8af481f26e050] --- .../include/hip/hcc_detail/hip_runtime_api.h | 71 +++++- .../include/hip/nvcc_detail/hip_runtime_api.h | 14 ++ projects/hip/src/hip_stream.cpp | 47 +++- .../stream/hipStreamCreateWithPriority.cpp | 220 ++++++++++++++++++ 4 files changed, 339 insertions(+), 13 deletions(-) create mode 100644 projects/hip/tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h index 573ae39af9..60609fd135 100644 --- a/projects/hip/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/hcc_detail/hip_runtime_api.h @@ -590,8 +590,6 @@ const char* hipGetErrorString(hipError_t hipError); * * The following Stream APIs are not (yet) supported in HIP: * - cudaStreamAttachMemAsync - * - cudaStreamCreateWithPriority - * - cudaStreamGetPriority */ @@ -609,7 +607,7 @@ const char* hipGetErrorString(hipError_t hipError); * * @return #hipSuccess, #hipErrorInvalidValue * - * @see hipStreamCreateWithFlags, hipStreamSynchronize, hipStreamWaitEvent, hipStreamDestroy + * @see hipStreamCreateWithFlags, hipStreamCreateWithPriority, hipStreamSynchronize, hipStreamWaitEvent, hipStreamDestroy */ hipError_t hipStreamCreate(hipStream_t* stream); @@ -628,12 +626,50 @@ hipError_t hipStreamCreate(hipStream_t* stream); * stream. See #hipStreamDefault, #hipStreamNonBlocking. * * - * @see hipStreamCreate, hipStreamSynchronize, hipStreamWaitEvent, hipStreamDestroy + * @see hipStreamCreate, hipStreamCreateWithPriority, hipStreamSynchronize, hipStreamWaitEvent, hipStreamDestroy */ hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags); +/** + * @brief Create an asynchronous stream with the specified priority. + * + * @param[in, out] stream Pointer to new stream + * @param[in ] flags to control stream creation. + * @param[in ] priority of the stream. Lower numbers represent higher priorities. + * @return #hipSuccess, #hipErrorInvalidValue + * + * Create a new asynchronous stream with the specified priority. @p stream returns an opaque handle + * that can be used to reference the newly created stream in subsequent hipStream* commands. The + * stream is allocated on the heap and will remain allocated even if the handle goes out-of-scope. + * To release the memory used by the stream, applicaiton must call hipStreamDestroy. Flags controls + * behavior of the stream. See #hipStreamDefault, #hipStreamNonBlocking. + * + * + * @see hipStreamCreate, hipStreamSynchronize, hipStreamWaitEvent, hipStreamDestroy + */ + +hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, int priority); + + +/** + * @brief Returns numerical values that correspond to the least and greatest stream priority. + * + * @param[in, out] leastPriority pointer in which value corresponding to least priority is returned. + * @param[in, out] greatestPriority pointer in which value corresponding to greatest priority is returned. + * + * Returns in *leastPriority and *greatestPriority the numerical values that correspond to the least + * and greatest stream priority respectively. Stream priorities follow a convention where lower numbers + * imply greater priorities. The range of meaningful stream priorities is given by + * [*greatestPriority, *leastPriority]. If the user attempts to create a stream with a priority value + * that is outside the the meaningful range as specified by this API, the priority is automatically + * clamped to within the valid range. + */ + +hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPriority); + + /** * @brief Destroys the specified stream. * @@ -649,7 +685,7 @@ hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags); * The queue may be destroyed while some commands are still inflight, or may wait for all commands * queued to the stream before destroying it. * - * @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamQuery, hipStreamWaitEvent, + * @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamCreateWithPriority, hipStreamQuery, hipStreamWaitEvent, * hipStreamSynchronize */ hipError_t hipStreamDestroy(hipStream_t stream); @@ -667,7 +703,7 @@ hipError_t hipStreamDestroy(hipStream_t stream); * host threads are sending work to the stream, the status may change immediately after the function * is called. It is typically used for debug. * - * @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamWaitEvent, hipStreamSynchronize, + * @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamCreateWithPriority, hipStreamWaitEvent, hipStreamSynchronize, * hipStreamDestroy */ hipError_t hipStreamQuery(hipStream_t stream); @@ -689,7 +725,7 @@ hipError_t hipStreamQuery(hipStream_t stream); * This command honors the hipDeviceLaunchBlocking flag, which controls whether the wait is active * or blocking. * - * @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamWaitEvent, hipStreamDestroy + * @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamCreateWithPriority, hipStreamWaitEvent, hipStreamDestroy * */ hipError_t hipStreamSynchronize(hipStream_t stream); @@ -712,7 +748,7 @@ hipError_t hipStreamSynchronize(hipStream_t stream); * does not impliciy wait for commands in the default stream to complete, even if the specified * stream is created with hipStreamNonBlocking = 0. * - * @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamSynchronize, hipStreamDestroy + * @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamCreateWithPriority, hipStreamSynchronize, hipStreamDestroy */ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags); @@ -732,6 +768,23 @@ hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int */ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags); + +/** + * @brief Query the priority of a stream. + * + * @param[in] stream stream to be queried + * @param[in,out] priority Pointer to an unsigned integer in which the stream's priority is returned + * @return #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidResourceHandle + * + * @returns #hipSuccess #hipErrorInvalidValue #hipErrorInvalidResourceHandle + * + * Query the priority of a stream. The priority is returned in in priority. + * + * @see hipStreamCreateWithFlags + */ +hipError_t hipStreamGetPriority(hipStream_t stream, int* priority); + + /** * Stream CallBack struct */ @@ -749,7 +802,7 @@ typedef void (*hipStreamCallback_t)(hipStream_t stream, hipError_t status, void* * @return #hipSuccess, #hipErrorInvalidResourceHandle, #hipErrorNotSupported * * @see hipStreamCreate, hipStreamCreateWithFlags, hipStreamQuery, hipStreamSynchronize, - * hipStreamWaitEvent, hipStreamDestroy + * hipStreamWaitEvent, hipStreamDestroy, hipStreamCreateWithPriority * */ hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData, diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index 55a3794846..e8b2f7d317 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -890,6 +890,13 @@ inline static hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned return hipCUDAErrorTohipError(cudaStreamCreateWithFlags(stream, flags)); } +inline static hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, int priority) { + return hipCUDAErrorTohipError(cudaStreamCreateWithPriority(stream, flags, priority)); +} + +inline static hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPriority) { + return hipCUDAErrorTohipError(cudaDeviceGetStreamPriorityRange(leastPriority, greatestPriority)); +} inline static hipError_t hipStreamCreate(hipStream_t* stream) { return hipCUDAErrorTohipError(cudaStreamCreate(stream)); @@ -903,6 +910,13 @@ inline static hipError_t hipStreamDestroy(hipStream_t stream) { return hipCUDAErrorTohipError(cudaStreamDestroy(stream)); } +inline static hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int *flags) { + return hipCUDAErrorTohipError(cudaStreamGetFlags(stream, flags)); +} + +inline static hipError_t hipStreamGetPriority(hipStream_t stream, int *priority) { + return hipCUDAErrorTohipError(cudaStreamGetPriority(stream, priority)); +} inline static hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) { diff --git a/projects/hip/src/hip_stream.cpp b/projects/hip/src/hip_stream.cpp index 4170e33d63..1e239cb2fc 100644 --- a/projects/hip/src/hip_stream.cpp +++ b/projects/hip/src/hip_stream.cpp @@ -31,9 +31,15 @@ THE SOFTWARE. //------------------------------------------------------------------------------------------------- // Stream // +enum queue_priority +{ + priority_high = Kalmar::priority_high, + priority_normal = Kalmar::priority_normal, + priority_low = Kalmar::priority_low +}; //--- -hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags) { +hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priority) { ihipCtx_t* ctx = ihipGetTlsDefaultCtx(); hipError_t e = hipSuccess; @@ -53,7 +59,7 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags) { // Obtain mutex access to the device critical data, release by destructor LockedAccessor_CtxCrit_t ctxCrit(ctx->criticalData()); - auto istream = new ihipStream_t(ctx, acc.create_view(), flags); + auto istream = new ihipStream_t(ctx, acc.create_view(Kalmar::execute_in_order, Kalmar::queuing_mode_automatic, (Kalmar::queue_priority)priority), flags); ctxCrit->addStream(istream); *stream = istream; @@ -73,16 +79,33 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags) { hipError_t hipStreamCreateWithFlags(hipStream_t* stream, unsigned int flags) { HIP_INIT_API(stream, flags); - return ihipLogStatus(ihipStreamCreate(stream, flags)); + return ihipLogStatus(ihipStreamCreate(stream, flags, priority_normal)); } //--- hipError_t hipStreamCreate(hipStream_t* stream) { HIP_INIT_API(stream); - return ihipLogStatus(ihipStreamCreate(stream, hipStreamDefault)); + return ihipLogStatus(ihipStreamCreate(stream, hipStreamDefault, priority_normal)); } +//--- +hipError_t hipStreamCreateWithPriority(hipStream_t* stream, unsigned int flags, int priority) { + HIP_INIT_API(stream, flags, priority); + + // clamp priority to range [priority_high:priority_low] + priority = (priority < priority_high ? priority_high : (priority > priority_low ? priority_low : priority)); + return ihipLogStatus(ihipStreamCreate(stream, flags, priority)); +} + +//--- +hipError_t hipDeviceGetStreamPriorityRange(int* leastPriority, int* greatestPriority) { + HIP_INIT_API(leastPriority, greatestPriority); + + if (leastPriority != NULL) *leastPriority = priority_low; + if (greatestPriority != NULL) *greatestPriority = priority_high; + return ihipLogStatus(hipSuccess); +} hipError_t hipStreamWaitEvent(hipStream_t stream, hipEvent_t event, unsigned int flags) { HIP_INIT_SPECIAL_API(TRACE_SYNC, stream, event, flags); @@ -191,6 +214,22 @@ hipError_t hipStreamGetFlags(hipStream_t stream, unsigned int* flags) { } +//-- +hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) { + HIP_INIT_API(stream, priority); + + if (priority == NULL) { + return ihipLogStatus(hipErrorInvalidValue); + } else if (stream == hipStreamNull) { + return ihipLogStatus(hipErrorInvalidResourceHandle); + } else { + LockedAccessor_StreamCrit_t crit(stream->_criticalData); + *priority = crit->_av.get_queue_priority(); + return ihipLogStatus(hipSuccess); + } +} + + //--- hipError_t hipStreamAddCallback(hipStream_t stream, hipStreamCallback_t callback, void* userData, unsigned int flags) { diff --git a/projects/hip/tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp b/projects/hip/tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp new file mode 100644 index 0000000000..582da67d85 --- /dev/null +++ b/projects/hip/tests/src/runtimeApi/stream/hipStreamCreateWithPriority.cpp @@ -0,0 +1,220 @@ +/* +Copyright (c) 2015-2016 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 + */ + +#include "hip/hip_runtime.h" +#include "test_common.h" + +#define MEMCPYSIZE 64*1024*1024 +#define NUMITERS 2 +#define GRIDSIZE 1024 +#define BLOCKSIZE 256 + +// helper rountine to initialize memory +template +void mem_init(T* buf, size_t n) +{ + for (int i = 0; i < n; i++) + { + buf[i] = i; + } +} + +// kernel to copy n elements from src to dst +template +__global__ void memcpy_kernel(T* dst, T* src, size_t n) +{ + int num = gridDim.x * blockDim.x; + int id = blockDim.x * blockIdx.x + threadIdx.x; + + for (int i = id; i < n; i += num) + { + dst[i] = src[i]; + } +} + +template +void runTest() +{ + size_t size = NUMITERS*MEMCPYSIZE; + + // get the range of priorities available + #define OP(x) \ + int priority_##x; \ + bool enable_priority_##x = false; + OP(low) + OP(normal) + OP(high) + #undef OP + HIPCHECK(hipDeviceGetStreamPriorityRange(&priority_low, &priority_high)); + printf("HIP stream priority range - low: %d to high: %d\n", priority_low, priority_high); + + // Check if priorities are indeed supported + if ((priority_low - priority_high) == 0) { passed(); } + + // Enable/disable priorities based on number of available priority levels + enable_priority_low = true; + enable_priority_high = true; + if ((priority_low - priority_high) > 1) enable_priority_normal = true; + if (enable_priority_normal) priority_normal = ((priority_low - priority_high) / 2); + + // create streams with highest and lowest available priorities + #define OP(x) \ + hipStream_t stream_##x; \ + if (enable_priority_##x) { \ + HIPCHECK(hipStreamCreateWithPriority(&stream_##x, hipStreamDefault, priority_##x)); \ + } + OP(low) + OP(normal) + OP(high) + #undef OP + + // allocate and initialise host source and destination buffers + #define OP(x) \ + T* src_h_##x; \ + T* dst_h_##x; \ + if (enable_priority_##x) { \ + src_h_##x = (T*)malloc(size); \ + if (src_h_##x == NULL) { printf("src_h_%s malloc failed!\n", #x); exit(-1); } \ + mem_init(src_h_##x, (size / sizeof(T))); \ + dst_h_##x = (T*)malloc(size); \ + if (dst_h_##x == NULL) { printf("dst_h_%s malloc failed!\n", #x); exit(-1); } \ + memset(dst_h_##x, 0, size); \ + } + OP(low) + OP(normal) + OP(high) + #undef OP + + // allocate and initialize device source and destination buffers + #define OP(x) \ + T* src_d_##x; \ + T* dst_d_##x; \ + if (enable_priority_##x) { \ + HIPCHECK(hipMalloc(&src_d_##x, size)); \ + HIPCHECK(hipMemcpy(src_d_##x, src_h_##x, size, hipMemcpyHostToDevice)); \ + HIPCHECK(hipMalloc(&dst_d_##x, size)); \ + } + OP(low) + OP(normal) + OP(high) + #undef OP + + // create events for measuring time spent in kernel execution + #define OP(x) \ + hipEvent_t event_start_##x; \ + hipEvent_t event_end_##x; \ + if (enable_priority_##x) { \ + HIPCHECK(hipEventCreate(&event_start_##x)); \ + HIPCHECK(hipEventCreate(&event_end_##x)); \ + } + OP(low) + OP(normal) + OP(high) + #undef OP + + // record start events for each of the priority streams + #define OP(x) \ + if (enable_priority_##x) { \ + HIPCHECK(hipEventRecord(event_start_##x, stream_##x)); \ + } + OP(low) + OP(normal) + OP(high) + #undef OP + + // launch kernels repeatedly on each of the prioritiy streams + for (int i = 0; i < size; i += MEMCPYSIZE) + { + int j = i / sizeof(T); + #define OP(x) \ + if (enable_priority_##x) { \ + hipLaunchKernelGGL((memcpy_kernel), dim3(GRIDSIZE), dim3(BLOCKSIZE), 0, stream_##x, dst_d_##x + j, src_d_##x + j, (MEMCPYSIZE / sizeof(T))); \ + } + OP(low) + OP(normal) + OP(high) + #undef OP + } + + // record end events for each of the priority streams + #define OP(x) \ + if (enable_priority_##x) { \ + HIPCHECK(hipEventRecord(event_end_##x, stream_##x)); \ + } + OP(low) + OP(normal) + OP(high) + #undef OP + + // synchronize events for each of the priority streams + #define OP(x) \ + if (enable_priority_##x) { \ + HIPCHECK(hipEventSynchronize(event_end_##x)); \ + } + OP(low) + OP(normal) + OP(high) + #undef OP + + // compute time spent for memcpy in each stream + #define OP(x) \ + float time_spent_##x; \ + if (enable_priority_##x) { \ + HIPCHECK(hipEventElapsedTime(&time_spent_##x, event_start_##x, event_end_##x)); \ + printf("time spent for memcpy in %6s priority stream: %.3lf ms\n", #x, time_spent_##x); \ + } + OP(low) + OP(normal) + OP(high) + #undef OP + + // sanity check + #define OP(x) \ + if (enable_priority_##x) { \ + HIPCHECK(hipMemcpy(dst_h_##x, dst_d_##x, size, hipMemcpyDeviceToHost)); \ + if (memcmp(dst_h_##x, src_h_##x, size) != 0) { printf("memcmp for %s failed!\n", #x); exit(-1); } \ + } + OP(low) + OP(normal) + OP(high) + #undef OP + + // validate that stream priorities are working as expected + #define OP(x, y) \ + if (enable_priority_##x && enable_priority_##y) { \ + if (time_spent_##x < time_spent_##y) { printf("FAILED!"); exit(-1); } \ + } + OP(low, normal) + OP(normal, high) + OP(low, high) + #undef OP + passed(); +} + +int main(int argc, char **argv) +{ + HipTest::parseStandardArguments(argc, argv, false); + runTest(); +} From 83e56b2ed404a2ee87add4ac51db4ea16d64c62f Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 15 Oct 2018 15:30:24 +0530 Subject: [PATCH 6/9] Add missing hipHostRegister flags on nvcc path Change-Id: I69f09204d9c544935104d4168ab8d3626666a623 [ROCm/hip commit: 64d1cf86b786640eae41df54bcdf0ae8ff22d8e2] --- projects/hip/include/hip/nvcc_detail/hip_runtime_api.h | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h index 55a3794846..c729773e92 100644 --- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h @@ -105,8 +105,10 @@ typedef enum hipChannelFormatKind { #define hipHostMallocCoherent 0x0 #define hipHostMallocNonCoherent 0x0 +#define hipHostRegisterDefault cudaHostRegisterDefault #define hipHostRegisterPortable cudaHostRegisterPortable #define hipHostRegisterMapped cudaHostRegisterMapped +#define hipHostRegisterIoMemory cudaHostRegisterIoMemory #define HIP_LAUNCH_PARAM_BUFFER_POINTER CU_LAUNCH_PARAM_BUFFER_POINTER #define HIP_LAUNCH_PARAM_BUFFER_SIZE CU_LAUNCH_PARAM_BUFFER_SIZE From 5b192ed119e0c2812f3e61f44172561fd38488d6 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 15 Oct 2018 15:27:37 +0300 Subject: [PATCH 7/9] [HIPIFY] Code cleanup and formatting [ROCm/hip commit: ee3a5cc722fd582262b2c7d75cd49049680ff431] --- projects/hip/hipify-clang/src/ArgParse.h | 2 - projects/hip/hipify-clang/src/CUDA2HipMap.h | 21 +- .../hip/hipify-clang/src/HipifyAction.cpp | 708 +++++++++--------- projects/hip/hipify-clang/src/HipifyAction.h | 133 ++-- projects/hip/hipify-clang/src/LLVMCompat.cpp | 28 +- projects/hip/hipify-clang/src/LLVMCompat.h | 32 +- .../src/ReplacementsFrontendActionFactory.h | 25 +- projects/hip/hipify-clang/src/Statistics.cpp | 268 +++---- projects/hip/hipify-clang/src/Statistics.h | 252 +++---- projects/hip/hipify-clang/src/StringUtils.cpp | 18 +- projects/hip/hipify-clang/src/StringUtils.h | 8 +- projects/hip/hipify-clang/src/main.cpp | 32 +- 12 files changed, 684 insertions(+), 843 deletions(-) diff --git a/projects/hip/hipify-clang/src/ArgParse.h b/projects/hip/hipify-clang/src/ArgParse.h index b937a8dd15..609544b058 100644 --- a/projects/hip/hipify-clang/src/ArgParse.h +++ b/projects/hip/hipify-clang/src/ArgParse.h @@ -7,7 +7,6 @@ namespace cl = llvm::cl; namespace ct = clang::tooling; extern cl::OptionCategory ToolTemplateCategory; - extern cl::opt OutputFilename; extern cl::opt Inplace; extern cl::opt NoBackup; @@ -15,5 +14,4 @@ extern cl::opt NoOutput; extern cl::opt PrintStats; extern cl::opt OutputStatsFilename; extern cl::opt Examine; - extern cl::extrahelp CommonHelp; diff --git a/projects/hip/hipify-clang/src/CUDA2HipMap.h b/projects/hip/hipify-clang/src/CUDA2HipMap.h index 605acf7aac..f45ef00851 100644 --- a/projects/hip/hipify-clang/src/CUDA2HipMap.h +++ b/projects/hip/hipify-clang/src/CUDA2HipMap.h @@ -3,26 +3,25 @@ #include "llvm/ADT/StringRef.h" #include #include - #include "Statistics.h" #define HIP_UNSUPPORTED true -/// Maps cuda header names to hip header names. +// Maps cuda header names to hip header names. extern const std::map CUDA_INCLUDE_MAP; -/// Maps the names of CUDA types to the corresponding hip types. +// Maps the names of CUDA types to the corresponding hip types. extern const std::map CUDA_TYPE_NAME_MAP; -/// Map all other CUDA identifiers (function/macro names, enum values) to hip versions. +// Map all other CUDA identifiers (function/macro names, enum values) to hip versions. extern const std::map CUDA_IDENTIFIER_MAP; /** - * The union of all the above maps. - * - * This should be used rarely, but is still needed to convert macro definitions (which can - * contain any combination of the above things). AST walkers can usually get away with just - * looking in the lookup table for the type of element they are processing, however, saving - * a great deal of time. - */ + * The union of all the above maps, except includes. + * + * This should be used rarely, but is still needed to convert macro definitions (which can + * contain any combination of the above things). AST walkers can usually get away with just + * looking in the lookup table for the type of element they are processing, however, saving + * a great deal of time. + */ const std::map& CUDA_RENAMES_MAP(); diff --git a/projects/hip/hipify-clang/src/HipifyAction.cpp b/projects/hip/hipify-clang/src/HipifyAction.cpp index a9bd1aa085..dab2cbf160 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.cpp +++ b/projects/hip/hipify-clang/src/HipifyAction.cpp @@ -1,12 +1,8 @@ #include "HipifyAction.h" - -#include - #include "clang/Basic/SourceLocation.h" #include "clang/Frontend/CompilerInstance.h" #include "clang/ASTMatchers/ASTMatchFinder.h" #include "clang/ASTMatchers/ASTMatchers.h" - #include "LLVMCompat.h" #include "CUDA2HipMap.h" #include "StringUtils.h" @@ -16,180 +12,167 @@ namespace ct = clang::tooling; namespace mat = clang::ast_matchers; void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) { - clang::SourceManager& SM = getCompilerInstance().getSourceManager(); - - size_t begin = 0; - while ((begin = s.find("cu", begin)) != StringRef::npos) { - const size_t end = s.find_first_of(" ", begin + 4); - StringRef name = s.slice(begin, end); - const auto found = CUDA_RENAMES_MAP().find(name); - if (found != CUDA_RENAMES_MAP().end()) { - StringRef repName = found->second.hipName; - hipCounter counter = {"[string literal]", ConvTypes::CONV_LITERAL, ApiTypes::API_RUNTIME, found->second.unsupported}; - Statistics::current().incrementCounter(counter, name.str()); - - if (!counter.unsupported) { - clang::SourceLocation sl = start.getLocWithOffset(begin + 1); - ct::Replacement Rep(SM, sl, name.size(), repName); - clang::FullSourceLoc fullSL(sl, SM); - insertReplacement(Rep, fullSL); - } - } - - if (end == StringRef::npos) { - break; - } - - begin = end + 1; + clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + size_t begin = 0; + while ((begin = s.find("cu", begin)) != StringRef::npos) { + const size_t end = s.find_first_of(" ", begin + 4); + StringRef name = s.slice(begin, end); + const auto found = CUDA_RENAMES_MAP().find(name); + if (found != CUDA_RENAMES_MAP().end()) { + StringRef repName = found->second.hipName; + hipCounter counter = {"[string literal]", ConvTypes::CONV_LITERAL, ApiTypes::API_RUNTIME, found->second.unsupported}; + Statistics::current().incrementCounter(counter, name.str()); + if (!counter.unsupported) { + clang::SourceLocation sl = start.getLocWithOffset(begin + 1); + ct::Replacement Rep(SM, sl, name.size(), repName); + clang::FullSourceLoc fullSL(sl, SM); + insertReplacement(Rep, fullSL); + } } + if (end == StringRef::npos) { + break; + } + begin = end + 1; + } } /** - * Look at, and consider altering, a given token. - * - * If it's not a CUDA identifier, nothing happens. - * If it's an unsupported CUDA identifier, a warning is emitted. - * Otherwise, the source file is updated with the corresponding hipification. - */ + * Look at, and consider altering, a given token. + * + * If it's not a CUDA identifier, nothing happens. + * If it's an unsupported CUDA identifier, a warning is emitted. + * Otherwise, the source file is updated with the corresponding hipification. + */ void HipifyAction::RewriteToken(const clang::Token& t) { - clang::SourceManager& SM = getCompilerInstance().getSourceManager(); - - // String literals containing CUDA references need fixing... - if (t.is(clang::tok::string_literal)) { - StringRef s(t.getLiteralData(), t.getLength()); - RewriteString(unquoteStr(s), t.getLocation()); - return; - } else if (!t.isAnyIdentifier()) { - // If it's neither a string nor an identifier, we don't care. - return; - } - - StringRef name = t.getRawIdentifier(); - const auto found = CUDA_RENAMES_MAP().find(name); - if (found == CUDA_RENAMES_MAP().end()) { - // So it's an identifier, but not CUDA? Boring. - return; - } - - Statistics::current().incrementCounter(found->second, name.str()); - - clang::SourceLocation sl = t.getLocation(); - if (found->second.unsupported) { - // An unsupported identifier? Curses! Warn the user. - clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics(); - const auto ID = DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "CUDA identifier unsupported in hip"); - DE.Report(sl, ID); - return; - } - - StringRef repName = found->second.hipName; - ct::Replacement Rep(SM, sl, name.size(), repName); - clang::FullSourceLoc fullSL(sl, SM); - insertReplacement(Rep, fullSL); + clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + // String literals containing CUDA references need fixing... + if (t.is(clang::tok::string_literal)) { + StringRef s(t.getLiteralData(), t.getLength()); + RewriteString(unquoteStr(s), t.getLocation()); + return; + } else if (!t.isAnyIdentifier()) { + // If it's neither a string nor an identifier, we don't care. + return; + } + StringRef name = t.getRawIdentifier(); + const auto found = CUDA_RENAMES_MAP().find(name); + if (found == CUDA_RENAMES_MAP().end()) { + // So it's an identifier, but not CUDA? Boring. + return; + } + Statistics::current().incrementCounter(found->second, name.str()); + clang::SourceLocation sl = t.getLocation(); + if (found->second.unsupported) { + // Warn the user about unsupported identifier. + clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics(); + const auto ID = DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "CUDA identifier unsupported in hip"); + DE.Report(sl, ID); + return; + } + StringRef repName = found->second.hipName; + ct::Replacement Rep(SM, sl, name.size(), repName); + clang::FullSourceLoc fullSL(sl, SM); + insertReplacement(Rep, fullSL); } namespace { clang::SourceRange getReadRange(clang::SourceManager& SM, const clang::SourceRange& exprRange) { - clang::SourceLocation begin = exprRange.getBegin(); - clang::SourceLocation end = exprRange.getEnd(); + clang::SourceLocation begin = exprRange.getBegin(); + clang::SourceLocation end = exprRange.getEnd(); - bool beginSafe = !SM.isMacroBodyExpansion(begin) || clang::Lexer::isAtStartOfMacroExpansion(begin, SM, clang::LangOptions{}); - bool endSafe = !SM.isMacroBodyExpansion(end) || clang::Lexer::isAtEndOfMacroExpansion(end, SM, clang::LangOptions{}); + bool beginSafe = !SM.isMacroBodyExpansion(begin) || clang::Lexer::isAtStartOfMacroExpansion(begin, SM, clang::LangOptions{}); + bool endSafe = !SM.isMacroBodyExpansion(end) || clang::Lexer::isAtEndOfMacroExpansion(end, SM, clang::LangOptions{}); - if (beginSafe && endSafe) { - return {SM.getFileLoc(begin), SM.getFileLoc(end)}; - } else { - return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)}; - } + if (beginSafe && endSafe) { + return {SM.getFileLoc(begin), SM.getFileLoc(end)}; + } else { + return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)}; + } } - clang::SourceRange getWriteRange(clang::SourceManager& SM, const clang::SourceRange& exprRange) { - clang::SourceLocation begin = exprRange.getBegin(); - clang::SourceLocation end = exprRange.getEnd(); - - // If the range is contained within a macro, update the macro definition. - // Otherwise, use the file location and hope for the best. - if (!SM.isMacroBodyExpansion(begin) || !SM.isMacroBodyExpansion(end)) { - return {SM.getFileLoc(begin), SM.getFileLoc(end)}; - } - - return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)}; + clang::SourceLocation begin = exprRange.getBegin(); + clang::SourceLocation end = exprRange.getEnd(); + // If the range is contained within a macro, update the macro definition. + // Otherwise, use the file location and hope for the best. + if (!SM.isMacroBodyExpansion(begin) || !SM.isMacroBodyExpansion(end)) { + return {SM.getFileLoc(begin), SM.getFileLoc(end)}; + } + return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)}; } - StringRef readSourceText(clang::SourceManager& SM, const clang::SourceRange& exprRange) { - return clang::Lexer::getSourceText(clang::CharSourceRange::getTokenRange(getReadRange(SM, exprRange)), SM, clang::LangOptions(), nullptr); + return clang::Lexer::getSourceText(clang::CharSourceRange::getTokenRange(getReadRange(SM, exprRange)), SM, clang::LangOptions(), nullptr); } /** - * Get a string representation of the expression `arg`, unless it's a defaulting function - * call argument, in which case get a 0. Used for building argument lists to kernel calls. - */ + * Get a string representation of the expression `arg`, unless it's a defaulting function + * call argument, in which case get a 0. Used for building argument lists to kernel calls. + */ std::string stringifyZeroDefaultedArg(clang::SourceManager& SM, const clang::Expr* arg) { - if (clang::isa(arg)) { - return "0"; - } else { - return readSourceText(SM, arg->getSourceRange()); - } + if (clang::isa(arg)) { + return "0"; + } else { + return readSourceText(SM, arg->getSourceRange()); + } } } // anonymous namespace bool HipifyAction::Exclude(const hipCounter & hipToken) { - switch (hipToken.type) { - case CONV_INCLUDE_CUDA_MAIN_H: - switch (hipToken.apiType) { - case API_DRIVER: - case API_RUNTIME: - if (insertedRuntimeHeader) { return true; } - insertedRuntimeHeader = true; - return false; - case API_BLAS: - if (insertedBLASHeader) { return true; } - insertedBLASHeader = true; - return false; - case API_RAND: - if (hipToken.hipName == "hiprand_kernel.h") { - if (insertedRAND_kernelHeader) { return true; } - insertedRAND_kernelHeader = true; - return false; - } else if (hipToken.hipName == "hiprand.h") { - if (insertedRANDHeader) { return true; } - insertedRANDHeader = true; - return false; - } - case API_DNN: - if (insertedDNNHeader) { return true; } - insertedDNNHeader = true; - return false; - case API_FFT: - if (insertedFFTHeader) { return true; } - insertedFFTHeader = true; - return false; - case API_COMPLEX: - if (insertedComplexHeader) { return true; } - insertedComplexHeader = true; - return false; - default: - return false; - } + switch (hipToken.type) { + case CONV_INCLUDE_CUDA_MAIN_H: + switch (hipToken.apiType) { + case API_DRIVER: + case API_RUNTIME: + if (insertedRuntimeHeader) { return true; } + insertedRuntimeHeader = true; + return false; + case API_BLAS: + if (insertedBLASHeader) { return true; } + insertedBLASHeader = true; + return false; + case API_RAND: + if (hipToken.hipName == "hiprand_kernel.h") { + if (insertedRAND_kernelHeader) { return true; } + insertedRAND_kernelHeader = true; return false; - case CONV_INCLUDE: - switch (hipToken.apiType) { - case API_RAND: - if (insertedRAND_kernelHeader) { return true; } - insertedRAND_kernelHeader = true; - return false; - default: - return false; - } + } else if (hipToken.hipName == "hiprand.h") { + if (insertedRANDHeader) { return true; } + insertedRANDHeader = true; return false; + } + case API_DNN: + if (insertedDNNHeader) { return true; } + insertedDNNHeader = true; + return false; + case API_FFT: + if (insertedFFTHeader) { return true; } + insertedFFTHeader = true; + return false; + case API_COMPLEX: + if (insertedComplexHeader) { return true; } + insertedComplexHeader = true; + return false; default: - return false; - } - return false; + return false; + } + return false; + case CONV_INCLUDE: + switch (hipToken.apiType) { + case API_RAND: + if (insertedRAND_kernelHeader) { return true; } + insertedRAND_kernelHeader = true; + return false; + default: + return false; + } + return false; + default: + return false; + } + return false; } void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, @@ -199,286 +182,263 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, clang::CharSourceRange filename_range, const clang::FileEntry*, StringRef, StringRef, const clang::Module*) { - clang::SourceManager& SM = getCompilerInstance().getSourceManager(); - if (!SM.isWrittenInMainFile(hash_loc)) { - return; - } - if (!firstHeader) { - firstHeader = true; - firstHeaderLoc = hash_loc; - } + clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + if (!SM.isWrittenInMainFile(hash_loc)) { + return; + } + if (!firstHeader) { + firstHeader = true; + firstHeaderLoc = hash_loc; + } + const auto found = CUDA_INCLUDE_MAP.find(file_name); + if (found == CUDA_INCLUDE_MAP.end()) { + return; + } + bool exclude = Exclude(found->second); + Statistics::current().incrementCounter(found->second, file_name.str()); - const auto found = CUDA_INCLUDE_MAP.find(file_name); - if (found == CUDA_INCLUDE_MAP.end()) { - return; - } + clang::SourceLocation sl = filename_range.getBegin(); + if (found->second.unsupported) { + clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics(); + DE.Report(sl, DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "Unsupported CUDA header")); + return; + } - bool exclude = Exclude(found->second); - - Statistics::current().incrementCounter(found->second, file_name.str()); - - clang::SourceLocation sl = filename_range.getBegin(); - if (found->second.unsupported) { - clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics(); - DE.Report(sl, DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "Unsupported CUDA header")); - return; - } - - clang::StringRef newInclude; - - // Keep the same include type that the user gave. - if (!exclude) { - clang::SmallString<128> includeBuffer; - if (is_angled) { - newInclude = llvm::Twine("<" + found->second.hipName + ">").toStringRef(includeBuffer); - } else { - newInclude = llvm::Twine("\"" + found->second.hipName + "\"").toStringRef(includeBuffer); - } + clang::StringRef newInclude; + // Keep the same include type that the user gave. + if (!exclude) { + clang::SmallString<128> includeBuffer; + if (is_angled) { + newInclude = llvm::Twine("<" + found->second.hipName + ">").toStringRef(includeBuffer); } else { - // hashLoc is location of the '#', thus replacing the whole include directive by empty newInclude starting with '#'. - sl = hash_loc; + newInclude = llvm::Twine("\"" + found->second.hipName + "\"").toStringRef(includeBuffer); } - const char *B = SM.getCharacterData(sl); - const char *E = SM.getCharacterData(filename_range.getEnd()); - ct::Replacement Rep(SM, sl, E - B, newInclude); - insertReplacement(Rep, clang::FullSourceLoc{sl, SM}); + } else { + // hashLoc is location of the '#', thus replacing the whole include directive by empty newInclude starting with '#'. + sl = hash_loc; + } + const char *B = SM.getCharacterData(sl); + const char *E = SM.getCharacterData(filename_range.getEnd()); + ct::Replacement Rep(SM, sl, E - B, newInclude); + insertReplacement(Rep, clang::FullSourceLoc{sl, SM}); } void HipifyAction::PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer) { - if (pragmaOnce) { - return; - } - clang::SourceManager& SM = getCompilerInstance().getSourceManager(); - if (!SM.isWrittenInMainFile(Loc)) { - return; - } - clang::Preprocessor& PP = getCompilerInstance().getPreprocessor(); - const clang::Token tok = PP.LookAhead(0); - StringRef Text(SM.getCharacterData(tok.getLocation()), tok.getLength()); - if (Text == "once") { - pragmaOnce = true; - pragmaOnceLoc = PP.LookAhead(1).getLocation(); - } + if (pragmaOnce) { + return; + } + clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + if (!SM.isWrittenInMainFile(Loc)) { + return; + } + clang::Preprocessor& PP = getCompilerInstance().getPreprocessor(); + const clang::Token tok = PP.LookAhead(0); + StringRef Text(SM.getCharacterData(tok.getLocation()), tok.getLength()); + if (Text == "once") { + pragmaOnce = true; + pragmaOnceLoc = PP.LookAhead(1).getLocation(); + } } bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::MatchResult& Result) { - StringRef refName = "cudaLaunchKernel"; + StringRef refName = "cudaLaunchKernel"; + const auto* launchKernel = Result.Nodes.getNodeAs(refName); + if (!launchKernel) { + return false; + } + clang::SmallString<40> XStr; + llvm::raw_svector_ostream OS(XStr); + clang::LangOptions DefaultLangOptions; + clang::SourceManager* SM = Result.SourceManager; - const auto* launchKernel = Result.Nodes.getNodeAs(refName); - if (!launchKernel) { - return false; - } + const clang::Expr& calleeExpr = *(launchKernel->getCallee()); + OS << "hipLaunchKernelGGL(" << readSourceText(*SM, calleeExpr.getSourceRange()) << ", "; - clang::SmallString<40> XStr; - llvm::raw_svector_ostream OS(XStr); + // Next up are the four kernel configuration parameters, the last two of which are optional and default to zero. + const clang::CallExpr& config = *(launchKernel->getConfig()); - clang::LangOptions DefaultLangOptions; - clang::SourceManager* SM = Result.SourceManager; + // Copy the two dimensional arguments verbatim. + OS << "dim3(" << readSourceText(*SM, config.getArg(0)->getSourceRange()) << "), "; + OS << "dim3(" << readSourceText(*SM, config.getArg(1)->getSourceRange()) << "), "; - const clang::Expr& calleeExpr = *(launchKernel->getCallee()); - OS << "hipLaunchKernelGGL(" << readSourceText(*SM, calleeExpr.getSourceRange()) << ", "; + // The stream/memory arguments default to zero if omitted. + OS << stringifyZeroDefaultedArg(*SM, config.getArg(2)) << ", "; + OS << stringifyZeroDefaultedArg(*SM, config.getArg(3)); - // Next up are the four kernel configuration parameters, the last two of which are optional and default to zero. - const clang::CallExpr& config = *(launchKernel->getConfig()); + // If there are ordinary arguments to the kernel, just copy them verbatim into our new call. + int numArgs = launchKernel->getNumArgs(); + if (numArgs > 0) { + OS << ", "; + // Start of the first argument. + clang::SourceLocation argStart = launchKernel->getArg(0)->getLocStart(); + // End of the last argument. + clang::SourceLocation argEnd = launchKernel->getArg(numArgs - 1)->getLocEnd(); + OS << readSourceText(*SM, {argStart, argEnd}); + } + OS << ")"; - // Copy the two dimensional arguments verbatim. - OS << "dim3(" << readSourceText(*SM, config.getArg(0)->getSourceRange()) << "), "; - OS << "dim3(" << readSourceText(*SM, config.getArg(1)->getSourceRange()) << "), "; - - // The stream/memory arguments default to zero if omitted. - OS << stringifyZeroDefaultedArg(*SM, config.getArg(2)) << ", "; - OS << stringifyZeroDefaultedArg(*SM, config.getArg(3)); - - // If there are ordinary arguments to the kernel, just copy them verbatim into our new call. - int numArgs = launchKernel->getNumArgs(); - if (numArgs > 0) { - OS << ", "; - - // Start of the first argument. - clang::SourceLocation argStart = launchKernel->getArg(0)->getLocStart(); - - // End of the last argument. - clang::SourceLocation argEnd = launchKernel->getArg(numArgs - 1)->getLocEnd(); - - OS << readSourceText(*SM, {argStart, argEnd}); - } - - OS << ")"; - - clang::SourceRange replacementRange = getWriteRange(*SM, {launchKernel->getLocStart(), launchKernel->getLocEnd()}); - clang::SourceLocation launchStart = replacementRange.getBegin(); - clang::SourceLocation launchEnd = replacementRange.getEnd(); - - size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchEnd, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(launchStart); - - ct::Replacement Rep(*SM, launchStart, length, OS.str()); - clang::FullSourceLoc fullSL(launchStart, *SM); - insertReplacement(Rep, fullSL); - hipCounter counter = {"hipLaunchKernelGGL", ConvTypes::CONV_KERN, ApiTypes::API_RUNTIME}; - Statistics::current().incrementCounter(counter, refName.str()); - - return true; + clang::SourceRange replacementRange = getWriteRange(*SM, {launchKernel->getLocStart(), launchKernel->getLocEnd()}); + clang::SourceLocation launchStart = replacementRange.getBegin(); + clang::SourceLocation launchEnd = replacementRange.getEnd(); + size_t length = SM->getCharacterData(clang::Lexer::getLocForEndOfToken(launchEnd, 0, *SM, DefaultLangOptions)) - SM->getCharacterData(launchStart); + ct::Replacement Rep(*SM, launchStart, length, OS.str()); + clang::FullSourceLoc fullSL(launchStart, *SM); + insertReplacement(Rep, fullSL); + hipCounter counter = {"hipLaunchKernelGGL", ConvTypes::CONV_KERN, ApiTypes::API_RUNTIME}; + Statistics::current().incrementCounter(counter, refName.str()); + return true; } bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::MatchFinder::MatchResult& Result) { - StringRef refName = "cudaSharedIncompleteArrayVar"; - auto* sharedVar = Result.Nodes.getNodeAs(refName); - if (!sharedVar) { - return false; - } + StringRef refName = "cudaSharedIncompleteArrayVar"; + auto* sharedVar = Result.Nodes.getNodeAs(refName); + if (!sharedVar) { + return false; + } + // Example: extern __shared__ uint sRadix1[]; + if (!sharedVar->hasExternalFormalLinkage()) { + return false; + } - // Example: extern __shared__ uint sRadix1[]; - if (!sharedVar->hasExternalFormalLinkage()) { - return false; + clang::QualType QT = sharedVar->getType(); + std::string typeName; + if (QT->isIncompleteArrayType()) { + const clang::ArrayType* AT = QT.getTypePtr()->getAsArrayTypeUnsafe(); + QT = AT->getElementType(); + if (QT.getTypePtr()->isBuiltinType()) { + QT = QT.getCanonicalType(); + const auto* BT = clang::dyn_cast(QT); + if (BT) { + clang::LangOptions LO; + LO.CUDA = true; + clang::PrintingPolicy policy(LO); + typeName = BT->getName(policy); + } + } else { + typeName = QT.getAsString(); } + } - clang::QualType QT = sharedVar->getType(); - std::string typeName; - if (QT->isIncompleteArrayType()) { - const clang::ArrayType* AT = QT.getTypePtr()->getAsArrayTypeUnsafe(); - QT = AT->getElementType(); - if (QT.getTypePtr()->isBuiltinType()) { - QT = QT.getCanonicalType(); - const auto* BT = clang::dyn_cast(QT); - if (BT) { - clang::LangOptions LO; - LO.CUDA = true; - clang::PrintingPolicy policy(LO); - typeName = BT->getName(policy); - } - } else { - typeName = QT.getAsString(); - } - } - - if (!typeName.empty()) { - clang::SourceLocation slStart = sharedVar->getLocStart(); - clang::SourceLocation slEnd = sharedVar->getLocEnd(); - clang::SourceManager* SM = Result.SourceManager; - size_t repLength = SM->getCharacterData(slEnd) - SM->getCharacterData(slStart) + 1; - std::string varName = sharedVar->getNameAsString(); - std::string repName = "HIP_DYNAMIC_SHARED(" + typeName + ", " + varName + ")"; - ct::Replacement Rep(*SM, slStart, repLength, repName); - clang::FullSourceLoc fullSL(slStart, *SM); - insertReplacement(Rep, fullSL); - hipCounter counter = {"HIP_DYNAMIC_SHARED", ConvTypes::CONV_MEM, ApiTypes::API_RUNTIME}; - Statistics::current().incrementCounter(counter, refName.str()); - } - - return true; + if (!typeName.empty()) { + clang::SourceLocation slStart = sharedVar->getLocStart(); + clang::SourceLocation slEnd = sharedVar->getLocEnd(); + clang::SourceManager* SM = Result.SourceManager; + size_t repLength = SM->getCharacterData(slEnd) - SM->getCharacterData(slStart) + 1; + std::string varName = sharedVar->getNameAsString(); + std::string repName = "HIP_DYNAMIC_SHARED(" + typeName + ", " + varName + ")"; + ct::Replacement Rep(*SM, slStart, repLength, repName); + clang::FullSourceLoc fullSL(slStart, *SM); + insertReplacement(Rep, fullSL); + hipCounter counter = {"HIP_DYNAMIC_SHARED", ConvTypes::CONV_MEM, ApiTypes::API_RUNTIME}; + Statistics::current().incrementCounter(counter, refName.str()); + } + return true; } void HipifyAction::insertReplacement(const ct::Replacement& rep, const clang::FullSourceLoc& fullSL) { - llcompat::insertReplacement(*replacements, rep); - if (PrintStats) { - rep.getLength(); - Statistics::current().lineTouched(fullSL.getExpansionLineNumber()); - Statistics::current().bytesChanged(rep.getLength()); - } + llcompat::insertReplacement(*replacements, rep); + if (PrintStats) { + rep.getLength(); + Statistics::current().lineTouched(fullSL.getExpansionLineNumber()); + Statistics::current().bytesChanged(rep.getLength()); + } } std::unique_ptr HipifyAction::CreateASTConsumer(clang::CompilerInstance& CI, llvm::StringRef) { - Finder.reset(new clang::ast_matchers::MatchFinder); - - // Replace the <<<...>>> language extension with a hip kernel launch - Finder->addMatcher(mat::cudaKernelCallExpr(mat::isExpansionInMainFile()).bind("cudaLaunchKernel"), this); - - Finder->addMatcher( - mat::varDecl( - mat::isExpansionInMainFile(), - mat::allOf( - mat::hasAttr(clang::attr::CUDAShared), - mat::hasType(mat::incompleteArrayType()) - ) - ).bind("cudaSharedIncompleteArrayVar"), - this - ); - - // Ownership is transferred to the caller... - return Finder->newASTConsumer(); + Finder.reset(new clang::ast_matchers::MatchFinder); + // Replace the <<<...>>> language extension with a hip kernel launch + Finder->addMatcher(mat::cudaKernelCallExpr(mat::isExpansionInMainFile()).bind("cudaLaunchKernel"), this); + Finder->addMatcher( + mat::varDecl( + mat::isExpansionInMainFile(), + mat::allOf( + mat::hasAttr(clang::attr::CUDAShared), + mat::hasType(mat::incompleteArrayType()) + ) + ).bind("cudaSharedIncompleteArrayVar"), + this + ); + // Ownership is transferred to the caller... + return Finder->newASTConsumer(); } void HipifyAction::EndSourceFileAction() { - // Insert the hip header, if we didn't already do it by accident during substitution. - if (!insertedRuntimeHeader) { - // It's not sufficient to just replace CUDA headers with hip ones, because numerous CUDA headers are - // implicitly included by the compiler. Instead, we _delete_ CUDA headers, and unconditionally insert - // one copy of the hip include into every file. - clang::SourceManager& SM = getCompilerInstance().getSourceManager(); - clang::SourceLocation sl; - if (pragmaOnce) { - sl = pragmaOnceLoc; - } else if (firstHeader) { - sl = firstHeaderLoc; - } else { - sl = SM.getLocForStartOfFile(SM.getMainFileID()); - } - clang::FullSourceLoc fullSL(sl, SM); - ct::Replacement Rep(SM, sl, 0, "\n#include \n"); - insertReplacement(Rep, fullSL); + // Insert the hip header, if we didn't already do it by accident during substitution. + if (!insertedRuntimeHeader) { + // It's not sufficient to just replace CUDA headers with hip ones, because numerous CUDA headers are + // implicitly included by the compiler. Instead, we _delete_ CUDA headers, and unconditionally insert + // one copy of the hip include into every file. + clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + clang::SourceLocation sl; + if (pragmaOnce) { + sl = pragmaOnceLoc; + } else if (firstHeader) { + sl = firstHeaderLoc; + } else { + sl = SM.getLocForStartOfFile(SM.getMainFileID()); } - - clang::ASTFrontendAction::EndSourceFileAction(); + clang::FullSourceLoc fullSL(sl, SM); + ct::Replacement Rep(SM, sl, 0, "\n#include \n"); + insertReplacement(Rep, fullSL); + } + clang::ASTFrontendAction::EndSourceFileAction(); } - namespace { /** - * A silly little class to proxy PPCallbacks back to the HipifyAction class. - */ + * A silly little class to proxy PPCallbacks back to the HipifyAction class. + */ class PPCallbackProxy : public clang::PPCallbacks { - HipifyAction& hipifyAction; + HipifyAction& hipifyAction; public: - explicit PPCallbackProxy(HipifyAction& action): hipifyAction(action) {} + explicit PPCallbackProxy(HipifyAction& action): hipifyAction(action) {} - void InclusionDirective(clang::SourceLocation hash_loc, const clang::Token& include_token, - StringRef file_name, bool is_angled, clang::CharSourceRange filename_range, - const clang::FileEntry* file, StringRef search_path, StringRef relative_path, - const clang::Module* imported + void InclusionDirective(clang::SourceLocation hash_loc, const clang::Token& include_token, + StringRef file_name, bool is_angled, clang::CharSourceRange filename_range, + const clang::FileEntry* file, StringRef search_path, StringRef relative_path, + const clang::Module* imported #if LLVM_VERSION_MAJOR > 6 - , clang::SrcMgr::CharacteristicKind FileType + , clang::SrcMgr::CharacteristicKind FileType #endif - ) override { - hipifyAction.InclusionDirective(hash_loc, include_token, file_name, is_angled, filename_range, file, search_path, relative_path, imported); - } + ) override { + hipifyAction.InclusionDirective(hash_loc, include_token, file_name, is_angled, filename_range, file, search_path, relative_path, imported); + } - void PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer) override { - hipifyAction.PragmaDirective(Loc, Introducer); - } + void PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer) override { + hipifyAction.PragmaDirective(Loc, Introducer); + } }; } void HipifyAction::ExecuteAction() { - clang::Preprocessor& PP = getCompilerInstance().getPreprocessor(); - clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + clang::Preprocessor& PP = getCompilerInstance().getPreprocessor(); + clang::SourceManager& SM = getCompilerInstance().getSourceManager(); - // Start lexing the specified input file. - const llvm::MemoryBuffer* FromFile = SM.getBuffer(SM.getMainFileID()); - clang::Lexer RawLex(SM.getMainFileID(), FromFile, SM, PP.getLangOpts()); - RawLex.SetKeepWhitespaceMode(true); + // Start lexing the specified input file. + const llvm::MemoryBuffer* FromFile = SM.getBuffer(SM.getMainFileID()); + clang::Lexer RawLex(SM.getMainFileID(), FromFile, SM, PP.getLangOpts()); + RawLex.SetKeepWhitespaceMode(true); - // Perform a token-level rewrite of CUDA identifiers to hip ones. The raw-mode lexer gives us enough - // information to tell the difference between identifiers, string literals, and "other stuff". It also - // ignores preprocessor directives, so this transformation will operate inside preprocessor-deleted - // code. - clang::Token RawTok; + // Perform a token-level rewrite of CUDA identifiers to hip ones. The raw-mode lexer gives us enough + // information to tell the difference between identifiers, string literals, and "other stuff". It also + // ignores preprocessor directives, so this transformation will operate inside preprocessor-deleted code. + clang::Token RawTok; + RawLex.LexFromRawLexer(RawTok); + while (RawTok.isNot(clang::tok::eof)) { + RewriteToken(RawTok); RawLex.LexFromRawLexer(RawTok); - while (RawTok.isNot(clang::tok::eof)) { - RewriteToken(RawTok); - RawLex.LexFromRawLexer(RawTok); - } + } - // Register yourself as the preprocessor callback, by proxy. - PP.addPPCallbacks(std::unique_ptr(new PPCallbackProxy(*this))); - - // Now we're done futzing with the lexer, have the subclass proceeed with Sema and AST matching. - clang::ASTFrontendAction::ExecuteAction(); + // Register yourself as the preprocessor callback, by proxy. + PP.addPPCallbacks(std::unique_ptr(new PPCallbackProxy(*this))); + // Now we're done futzing with the lexer, have the subclass proceeed with Sema and AST matching. + clang::ASTFrontendAction::ExecuteAction(); } void HipifyAction::run(const clang::ast_matchers::MatchFinder::MatchResult& Result) { - if (cudaLaunchKernel(Result)) return; - if (cudaSharedIncompleteArrayVar(Result)) return; + if (cudaLaunchKernel(Result)) return; + if (cudaSharedIncompleteArrayVar(Result)) return; } diff --git a/projects/hip/hipify-clang/src/HipifyAction.h b/projects/hip/hipify-clang/src/HipifyAction.h index 1262142cfc..7b54dddf54 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.h +++ b/projects/hip/hipify-clang/src/HipifyAction.h @@ -2,8 +2,8 @@ #include "clang/Lex/PPCallbacks.h" #include "clang/Tooling/Tooling.h" -#include "clang/Frontend/FrontendAction.h" #include "clang/Tooling/Core/Replacement.h" +#include "clang/Frontend/FrontendAction.h" #include "clang/ASTMatchers/ASTMatchFinder.h" #include "ReplacementsFrontendActionFactory.h" #include "Statistics.h" @@ -11,91 +11,62 @@ namespace ct = clang::tooling; /** - * A FrontendAction that hipifies CUDA programs. - */ + * A FrontendAction that hipifies CUDA programs. + */ class HipifyAction : public clang::ASTFrontendAction, public clang::ast_matchers::MatchFinder::MatchCallback { private: - ct::Replacements* replacements; - std::unique_ptr Finder; - - /// CUDA implicitly adds its runtime header. We rewrite explicitly-provided CUDA includes with equivalent - // ones, and track - using this flag - if the result led to us including the hip runtime header. If it did - // not, we insert it at the top of the file when we finish processing it. - // This approach means we do the best it's possible to do w.r.t preserving the user's include order. - bool insertedRuntimeHeader = false; - bool insertedBLASHeader = false; - bool insertedRANDHeader = false; - bool insertedRAND_kernelHeader = false; - bool insertedDNNHeader = false; - bool insertedFFTHeader = false; - bool insertedComplexHeader = false; - bool firstHeader = false; - bool pragmaOnce = false; - clang::SourceLocation firstHeaderLoc; - clang::SourceLocation pragmaOnceLoc; - - /** - * Rewrite a string literal to refer to hip, not CUDA. - */ - void RewriteString(StringRef s, clang::SourceLocation start); - - /** - * Replace a CUDA identifier with the corresponding hip identifier, if applicable. - */ - void RewriteToken(const clang::Token &t); + ct::Replacements* replacements; + std::unique_ptr Finder; + // CUDA implicitly adds its runtime header. We rewrite explicitly-provided CUDA includes with equivalent + // ones, and track - using this flag - if the result led to us including the hip runtime header. If it did + // not, we insert it at the top of the file when we finish processing it. + // This approach means we do the best it's possible to do w.r.t preserving the user's include order. + bool insertedRuntimeHeader = false; + bool insertedBLASHeader = false; + bool insertedRANDHeader = false; + bool insertedRAND_kernelHeader = false; + bool insertedDNNHeader = false; + bool insertedFFTHeader = false; + bool insertedComplexHeader = false; + bool firstHeader = false; + bool pragmaOnce = false; + clang::SourceLocation firstHeaderLoc; + clang::SourceLocation pragmaOnceLoc; + // Rewrite a string literal to refer to hip, not CUDA. + void RewriteString(StringRef s, clang::SourceLocation start); + // Replace a CUDA identifier with the corresponding hip identifier, if applicable. + void RewriteToken(const clang::Token &t); public: - explicit HipifyAction(ct::Replacements *replacements): - clang::ASTFrontendAction(), - replacements(replacements) {} - - // MatchCallback listeners - bool cudaBuiltin(const clang::ast_matchers::MatchFinder::MatchResult& Result); - bool cudaLaunchKernel(const clang::ast_matchers::MatchFinder::MatchResult& Result); - bool cudaSharedIncompleteArrayVar(const clang::ast_matchers::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, - StringRef file_name, - bool is_angled, - clang::CharSourceRange filename_range, - const clang::FileEntry *file, - StringRef search_path, - StringRef relative_path, - const clang::Module *imported); - - /** - * Called by the preprocessor for each pragma directive during the non-raw lexing pass. - */ - void PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer); + explicit HipifyAction(ct::Replacements *replacements): clang::ASTFrontendAction(), + replacements(replacements) {} + // MatchCallback listeners + bool cudaBuiltin(const clang::ast_matchers::MatchFinder::MatchResult& Result); + bool cudaLaunchKernel(const clang::ast_matchers::MatchFinder::MatchResult& Result); + bool cudaSharedIncompleteArrayVar(const clang::ast_matchers::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, + StringRef file_name, + bool is_angled, + clang::CharSourceRange filename_range, + const clang::FileEntry *file, + StringRef search_path, + StringRef relative_path, + const clang::Module *imported); + // Called by the preprocessor for each pragma directive during the non-raw lexing pass. + void PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer); protected: - /** - * Add a Replacement for the current file. These will all be applied after executing the FrontendAction. - */ - void insertReplacement(const ct::Replacement& rep, const clang::FullSourceLoc& fullSL); - - /** - * FrontendAction entry point. - */ - void ExecuteAction() override; - - /** - * Called at the start of each new file to process. - */ - void EndSourceFileAction() override; - - /** - * MatchCallback API entry point. Called by the AST visitor while searching the AST for things we registered an - * interest for. - */ - void run(const clang::ast_matchers::MatchFinder::MatchResult& Result) override; - - std::unique_ptr CreateASTConsumer(clang::CompilerInstance &CI, llvm::StringRef InFile) override; - - bool Exclude(const hipCounter & hipToken); + // Add a Replacement for the current file. These will all be applied after executing the FrontendAction. + void insertReplacement(const ct::Replacement& rep, const clang::FullSourceLoc& fullSL); + // FrontendAction entry point. + void ExecuteAction() override; + // Called at the start of each new file to process. + void EndSourceFileAction() override; + // MatchCallback API entry point. Called by the AST visitor while searching the AST for things we registered an interest for. + void run(const clang::ast_matchers::MatchFinder::MatchResult& Result) override; + std::unique_ptr CreateASTConsumer(clang::CompilerInstance &CI, llvm::StringRef InFile) override; + bool Exclude(const hipCounter & hipToken); }; diff --git a/projects/hip/hipify-clang/src/LLVMCompat.cpp b/projects/hip/hipify-clang/src/LLVMCompat.cpp index 6b6dc18dd2..4ab62310d6 100644 --- a/projects/hip/hipify-clang/src/LLVMCompat.cpp +++ b/projects/hip/hipify-clang/src/LLVMCompat.cpp @@ -3,40 +3,40 @@ namespace llcompat { void PrintStackTraceOnErrorSignal() { - // The signature of PrintStackTraceOnErrorSignal changed in llvm 3.9. We don't support - // anything older than 3.8, so let's specifically detect the one old version we support. + // The signature of PrintStackTraceOnErrorSignal changed in llvm 3.9. We don't support + // anything older than 3.8, so let's specifically detect the one old version we support. #if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8) - llvm::sys::PrintStackTraceOnErrorSignal(); + llvm::sys::PrintStackTraceOnErrorSignal(); #else - llvm::sys::PrintStackTraceOnErrorSignal(clang::StringRef()); + llvm::sys::PrintStackTraceOnErrorSignal(clang::StringRef()); #endif } ct::Replacements& getReplacements(ct::RefactoringTool& Tool, clang::StringRef file) { #if LLVM_VERSION_MAJOR > 3 - // getReplacements() now returns a map from filename to Replacements - so create an entry - // for this source file and return a reference to it. - return Tool.getReplacements()[file]; + // getReplacements() now returns a map from filename to Replacements - so create an entry + // for this source file and return a reference to it. + return Tool.getReplacements()[file]; #else - return Tool.getReplacements(); + return Tool.getReplacements(); #endif } void insertReplacement(ct::Replacements& replacements, const ct::Replacement& rep) { #if LLVM_VERSION_MAJOR > 3 - // New clang added error checking to Replacements, and *insists* that you explicitly check it. - llvm::consumeError(replacements.add(rep)); + // New clang added error checking to Replacements, and *insists* that you explicitly check it. + llvm::consumeError(replacements.add(rep)); #else - // In older versions, it's literally an std::set - replacements.insert(rep); + // In older versions, it's literally an std::set + replacements.insert(rep); #endif } void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token *start, size_t len, bool DisableMacroExpansion) { #if (LLVM_VERSION_MAJOR == 3) && (LLVM_VERSION_MINOR == 8) - _pp.EnterTokenStream(start, len, false, DisableMacroExpansion); + _pp.EnterTokenStream(start, len, false, DisableMacroExpansion); #else - _pp.EnterTokenStream(clang::ArrayRef{start, len}, DisableMacroExpansion); + _pp.EnterTokenStream(clang::ArrayRef{start, len}, DisableMacroExpansion); #endif } diff --git a/projects/hip/hipify-clang/src/LLVMCompat.h b/projects/hip/hipify-clang/src/LLVMCompat.h index 72b6832012..9f82e36a1f 100644 --- a/projects/hip/hipify-clang/src/LLVMCompat.h +++ b/projects/hip/hipify-clang/src/LLVMCompat.h @@ -11,40 +11,38 @@ namespace ct = clang::tooling; // Things for papering over the differences between different LLVM versions. namespace llcompat { - - /** - * The getNumArgs function on macros was rather unhelpfully renamed in clang 4.0. Its semantics - * remain unchanged, so let's be slightly ugly about it here. :D - */ + * The getNumArgs function on macros was rather unhelpfully renamed in clang 4.0. Its semantics + * remain unchanged, so let's be slightly ugly about it here. :D + */ #if LLVM_VERSION_MAJOR > 4 - #define GET_NUM_ARGS() getNumParams() + #define GET_NUM_ARGS() getNumParams() #else - #define GET_NUM_ARGS() getNumArgs() + #define GET_NUM_ARGS() getNumArgs() #endif #if LLVM_VERSION_MAJOR < 7 - #define LLVM_DEBUG(X) DEBUG(X) + #define LLVM_DEBUG(X) DEBUG(X) #endif void PrintStackTraceOnErrorSignal(); /** - * Get the replacement map for a given filename in a RefactoringTool. - * - * Older LLVM versions don't actually support multiple filenames, so everything all gets - * smushed together. It is the caller's responsibility to cope with this. - */ + * Get the replacement map for a given filename in a RefactoringTool. + * + * Older LLVM versions don't actually support multiple filenames, so everything all gets + * smushed together. It is the caller's responsibility to cope with this. + */ ct::Replacements& getReplacements(ct::RefactoringTool& Tool, clang::StringRef file); /** - * Add a Replacement to a Replacements. - */ + * Add a Replacement to a Replacements. + */ void insertReplacement(ct::Replacements& replacements, const ct::Replacement& rep); /** - * Version-agnostic version of Preprocessor::EnterTokenStream(). - */ + * Version-agnostic version of Preprocessor::EnterTokenStream(). + */ void EnterPreprocessorTokenStream(clang::Preprocessor& _pp, const clang::Token *start, size_t len, diff --git a/projects/hip/hipify-clang/src/ReplacementsFrontendActionFactory.h b/projects/hip/hipify-clang/src/ReplacementsFrontendActionFactory.h index 7896635ef6..9e0decdeb9 100644 --- a/projects/hip/hipify-clang/src/ReplacementsFrontendActionFactory.h +++ b/projects/hip/hipify-clang/src/ReplacementsFrontendActionFactory.h @@ -6,23 +6,22 @@ namespace ct = clang::tooling; - /** - * A FrontendActionFactory that propagates a set of Replacements into the FrontendAction. - * This is necessary boilerplate for using a custom FrontendAction with a RefactoringTool. - * - * @tparam T The FrontendAction to create. - */ + * A FrontendActionFactory that propagates a set of Replacements into the FrontendAction. + * This is necessary boilerplate for using a custom FrontendAction with a RefactoringTool. + * + * @tparam T The FrontendAction to create. + */ template class ReplacementsFrontendActionFactory : public ct::FrontendActionFactory { - ct::Replacements* replacements; + ct::Replacements* replacements; public: - explicit ReplacementsFrontendActionFactory(ct::Replacements* r): - ct::FrontendActionFactory(), - replacements(r) {} + explicit ReplacementsFrontendActionFactory(ct::Replacements* r): + ct::FrontendActionFactory(), + replacements(r) {} - clang::FrontendAction* create() override { - return new T(replacements); - } + clang::FrontendAction* create() override { + return new T(replacements); + } }; diff --git a/projects/hip/hipify-clang/src/Statistics.cpp b/projects/hip/hipify-clang/src/Statistics.cpp index 9b70a793ca..c012bf1131 100644 --- a/projects/hip/hipify-clang/src/Statistics.cpp +++ b/projects/hip/hipify-clang/src/Statistics.cpp @@ -3,18 +3,17 @@ #include #include - const char *counterNames[NUM_CONV_TYPES] = { - "version", "init", "device", "mem", "kern", "coord_func", "math_func", "device_func", - "special_func", "stream", "event", "occupancy", "ctx", "peer", "module", - "cache", "exec", "external_resource_interop", "graph", "err", "def", "tex", "gl", "graphics", - "surface", "jit", "d3d9", "d3d10", "d3d11", "vdpau", "egl", "complex", - "thread", "other", "include", "include_cuda_main_header", "type", "literal", - "numeric_literal" + "version", "init", "device", "mem", "kern", "coord_func", "math_func", "device_func", + "special_func", "stream", "event", "occupancy", "ctx", "peer", "module", + "cache", "exec", "external_resource_interop", "graph", "err", "def", "tex", "gl", "graphics", + "surface", "jit", "d3d9", "d3d10", "d3d11", "vdpau", "egl", "complex", + "thread", "other", "include", "include_cuda_main_header", "type", "literal", + "numeric_literal" }; const char *apiNames[NUM_API_TYPES] = { - "CUDA Driver API", "CUDA RT API", "CUBLAS API", "CURAND API", "CUDNN API", "CUFFT API", "cuComplex API" + "CUDA Driver API", "CUDA RT API", "CUBLAS API", "CURAND API", "CUDNN API", "CUFFT API", "cuComplex API" }; namespace { @@ -24,203 +23,174 @@ void conditionalPrint(ST *stream1, ST2* stream2, const std::string& s1, const std::string& s2) { - if (stream1) { - *stream1 << s1; - } - - if (stream2) { - *stream2 << s2; - } + if (stream1) { + *stream1 << s1; + } + if (stream2) { + *stream2 << s2; + } } - -/** - * Print a named stat value to both the terminal and the CSV file. - */ +// Print a named stat value to both the terminal and the CSV file. template void printStat(std::ostream *csv, llvm::raw_ostream* printOut, const std::string &name, T value) { - if (printOut) { - *printOut << " " << name << ": " << value << "\n"; - } - - if (csv) { - *csv << name << ";" << value << "\n"; - } + if (printOut) { + *printOut << " " << name << ": " << value << "\n"; + } + if (csv) { + *csv << name << ";" << value << "\n"; + } } - } // Anonymous namespace void StatCounter::incrementCounter(const hipCounter& counter, std::string name) { - counters[name]++; - apiCounters[(int) counter.apiType]++; - convTypeCounters[(int) counter.type]++; + counters[name]++; + apiCounters[(int) counter.apiType]++; + convTypeCounters[(int) counter.type]++; } void StatCounter::add(const StatCounter& other) { - for (const auto& p : other.counters) { - counters[p.first] += p.second; - } - - for (int i = 0; i < NUM_API_TYPES; i++) { - apiCounters[i] += other.apiCounters[i]; - } - - for (int i = 0; i < NUM_CONV_TYPES; i++) { - convTypeCounters[i] += other.convTypeCounters[i]; - } + for (const auto& p : other.counters) { + counters[p.first] += p.second; + } + for (int i = 0; i < NUM_API_TYPES; i++) { + apiCounters[i] += other.apiCounters[i]; + } + for (int i = 0; i < NUM_CONV_TYPES; i++) { + convTypeCounters[i] += other.convTypeCounters[i]; + } } int StatCounter::getConvSum() { - int acc = 0; - for (const int& i : convTypeCounters) { - acc += i; - } - - return acc; + int acc = 0; + for (const int& i : convTypeCounters) { + acc += i; + } + return acc; } void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, std::string prefix) { - conditionalPrint(csv, printOut, "\nCUDA ref type;Count\n", "[HIPIFY] info: " + prefix + " refs by type:\n"); - for (int i = 0; i < NUM_CONV_TYPES; i++) { - if (convTypeCounters[i] > 0) { - printStat(csv, printOut, counterNames[i], convTypeCounters[i]); - } - } - - conditionalPrint(csv, printOut, "\nCUDA API;Count\n", "[HIPIFY] info: " + prefix + " refs by API:\n"); - for (int i = 0; i < NUM_API_TYPES; i++) { - printStat(csv, printOut, apiNames[i], apiCounters[i]); - } - - conditionalPrint(csv, printOut, "\nCUDA ref name;Count\n", "[HIPIFY] info: " + prefix + " refs by names:\n"); - for (const auto &it : counters) { - printStat(csv, printOut, it.first, it.second); + conditionalPrint(csv, printOut, "\nCUDA ref type;Count\n", "[HIPIFY] info: " + prefix + " refs by type:\n"); + for (int i = 0; i < NUM_CONV_TYPES; i++) { + if (convTypeCounters[i] > 0) { + printStat(csv, printOut, counterNames[i], convTypeCounters[i]); } + } + conditionalPrint(csv, printOut, "\nCUDA API;Count\n", "[HIPIFY] info: " + prefix + " refs by API:\n"); + for (int i = 0; i < NUM_API_TYPES; i++) { + printStat(csv, printOut, apiNames[i], apiCounters[i]); + } + conditionalPrint(csv, printOut, "\nCUDA ref name;Count\n", "[HIPIFY] info: " + prefix + " refs by names:\n"); + for (const auto &it : counters) { + printStat(csv, printOut, it.first, it.second); + } } - Statistics::Statistics(std::string name): fileName(name) { - // Compute the total bytes/lines in the input file. - std::ifstream src_file(name, std::ios::binary | std::ios::ate); - src_file.clear(); - src_file.seekg(0); - totalLines = (int) std::count(std::istreambuf_iterator(src_file), std::istreambuf_iterator(), '\n'); - totalBytes = (int) src_file.tellg(); - - // Mark the start time... - startTime = chr::steady_clock::now(); -}; - + // Compute the total bytes/lines in the input file. + std::ifstream src_file(name, std::ios::binary | std::ios::ate); + src_file.clear(); + src_file.seekg(0); + totalLines = (int) std::count(std::istreambuf_iterator(src_file), std::istreambuf_iterator(), '\n'); + totalBytes = (int) src_file.tellg(); + startTime = chr::steady_clock::now(); +} ///////// Counter update routines ////////// void Statistics::incrementCounter(const hipCounter &counter, std::string name) { - if (counter.unsupported) { - unsupported.incrementCounter(counter, name); - } else { - supported.incrementCounter(counter, name); - } + if (counter.unsupported) { + unsupported.incrementCounter(counter, name); + } else { + supported.incrementCounter(counter, name); + } } void Statistics::add(const Statistics &other) { - supported.add(other.supported); - unsupported.add(other.unsupported); - totalBytes += other.totalBytes; - totalLines += other.totalLines; - touchedBytes += other.touchedBytes; + supported.add(other.supported); + unsupported.add(other.unsupported); + totalBytes += other.totalBytes; + totalLines += other.totalLines; + touchedBytes += other.touchedBytes; } void Statistics::lineTouched(int lineNumber) { - touchedLines.insert(lineNumber); + touchedLines.insert(lineNumber); } void Statistics::bytesChanged(int bytes) { - touchedBytes += bytes; + touchedBytes += bytes; } void Statistics::markCompletion() { - completionTime = chr::steady_clock::now(); + completionTime = chr::steady_clock::now(); } - ///////// Output functions ////////// void Statistics::print(std::ostream* csv, llvm::raw_ostream* printOut, bool skipHeader) { - if (!skipHeader) { - std::string str = "file \'" + fileName + "\' statistics:\n"; - conditionalPrint(csv, printOut, "\n" + str, "\n[HIPIFY] info: " + str); - } - - size_t changedLines = touchedLines.size(); - - // Total number of (un)supported refs that were converted. - int supportedSum = supported.getConvSum(); - int unsupportedSum = unsupported.getConvSum(); - - printStat(csv, printOut, "CONVERTED refs count", supportedSum); - printStat(csv, printOut, "UNCONVERTED refs count", unsupportedSum); - printStat(csv, printOut, "CONVERSION %", 100 - std::lround(double(unsupportedSum * 100) / double(supportedSum + unsupportedSum))); - printStat(csv, printOut, "REPLACED bytes", touchedBytes); - printStat(csv, printOut, "TOTAL bytes", totalBytes); - printStat(csv, printOut, "CHANGED lines of code", changedLines); - printStat(csv, printOut, "TOTAL lines of code", totalLines); - - if (totalBytes > 0) { - printStat(csv, printOut, "CODE CHANGED (in bytes) %", std::lround(double(touchedBytes * 100) / double(totalBytes))); - } - - if (totalLines > 0) { - printStat(csv, printOut, "CODE CHANGED (in lines) %", std::lround(double(changedLines * 100) / double(totalLines))); - } - - typedef std::chrono::duration duration; - duration elapsed = completionTime - startTime; - std::stringstream stream; - stream << std::fixed << std::setprecision(2) << elapsed.count() / 1000; - printStat(csv, printOut, "TIME ELAPSED s", stream.str()); - - supported.print(csv, printOut, "CONVERTED"); - unsupported.print(csv, printOut, "UNCONVERTED"); + if (!skipHeader) { + std::string str = "file \'" + fileName + "\' statistics:\n"; + conditionalPrint(csv, printOut, "\n" + str, "\n[HIPIFY] info: " + str); + } + size_t changedLines = touchedLines.size(); + // Total number of (un)supported refs that were converted. + int supportedSum = supported.getConvSum(); + int unsupportedSum = unsupported.getConvSum(); + printStat(csv, printOut, "CONVERTED refs count", supportedSum); + printStat(csv, printOut, "UNCONVERTED refs count", unsupportedSum); + printStat(csv, printOut, "CONVERSION %", 100 - std::lround(double(unsupportedSum * 100) / double(supportedSum + unsupportedSum))); + printStat(csv, printOut, "REPLACED bytes", touchedBytes); + printStat(csv, printOut, "TOTAL bytes", totalBytes); + printStat(csv, printOut, "CHANGED lines of code", changedLines); + printStat(csv, printOut, "TOTAL lines of code", totalLines); + if (totalBytes > 0) { + printStat(csv, printOut, "CODE CHANGED (in bytes) %", std::lround(double(touchedBytes * 100) / double(totalBytes))); + } + if (totalLines > 0) { + printStat(csv, printOut, "CODE CHANGED (in lines) %", std::lround(double(changedLines * 100) / double(totalLines))); + } + typedef std::chrono::duration duration; + duration elapsed = completionTime - startTime; + std::stringstream stream; + stream << std::fixed << std::setprecision(2) << elapsed.count() / 1000; + printStat(csv, printOut, "TIME ELAPSED s", stream.str()); + supported.print(csv, printOut, "CONVERTED"); + unsupported.print(csv, printOut, "UNCONVERTED"); } void Statistics::printAggregate(std::ostream *csv, llvm::raw_ostream* printOut) { - Statistics globalStats = getAggregate(); - - conditionalPrint(csv, printOut, "\nTOTAL statistics:\n", "\n[HIPIFY] info: TOTAL statistics:\n"); - - // A file is considered "converted" if we made any changes to it. - int convertedFiles = 0; - for (const auto& p : stats) { - if (!p.second.touchedLines.empty()) { - convertedFiles++; - } + Statistics globalStats = getAggregate(); + conditionalPrint(csv, printOut, "\nTOTAL statistics:\n", "\n[HIPIFY] info: TOTAL statistics:\n"); + // A file is considered "converted" if we made any changes to it. + int convertedFiles = 0; + for (const auto& p : stats) { + if (!p.second.touchedLines.empty()) { + convertedFiles++; } - - printStat(csv, printOut, "CONVERTED files", convertedFiles); - printStat(csv, printOut, "PROCESSED files", stats.size()); - - globalStats.print(csv, printOut); + } + printStat(csv, printOut, "CONVERTED files", convertedFiles); + printStat(csv, printOut, "PROCESSED files", stats.size()); + globalStats.print(csv, printOut); } //// Static state management //// Statistics Statistics::getAggregate() { - Statistics globalStats("global"); - - for (const auto& p : stats) { - globalStats.add(p.second); - } - - return globalStats; + Statistics globalStats("global"); + for (const auto& p : stats) { + globalStats.add(p.second); + } + return globalStats; } Statistics& Statistics::current() { - assert(Statistics::currentStatistics); - return *Statistics::currentStatistics; + assert(Statistics::currentStatistics); + return *Statistics::currentStatistics; } void Statistics::setActive(std::string name) { - stats.emplace(std::make_pair(name, Statistics{name})); - Statistics::currentStatistics = &stats.at(name); + stats.emplace(std::make_pair(name, Statistics{name})); + Statistics::currentStatistics = &stats.at(name); } std::map Statistics::stats = {}; diff --git a/projects/hip/hipify-clang/src/Statistics.h b/projects/hip/hipify-clang/src/Statistics.h index b53bcdf6a0..af8d23fd75 100644 --- a/projects/hip/hipify-clang/src/Statistics.h +++ b/projects/hip/hipify-clang/src/Statistics.h @@ -3,66 +3,66 @@ #include #include #include -#include #include #include +#include #include namespace chr = std::chrono; enum ConvTypes { - CONV_VERSION = 0, - CONV_INIT, - CONV_DEVICE, - CONV_MEM, - CONV_KERN, - CONV_COORD_FUNC, - CONV_MATH_FUNC, - CONV_DEVICE_FUNC, - CONV_SPECIAL_FUNC, - CONV_STREAM, - CONV_EVENT, - CONV_OCCUPANCY, - CONV_CONTEXT, - CONV_PEER, - CONV_MODULE, - CONV_CACHE, - CONV_EXEC, - CONV_EXTERNAL_RES, - CONV_GRAPH, - CONV_ERROR, - CONV_DEF, - CONV_TEX, - CONV_GL, - CONV_GRAPHICS, - CONV_SURFACE, - CONV_JIT, - CONV_D3D9, - CONV_D3D10, - CONV_D3D11, - CONV_VDPAU, - CONV_EGL, - CONV_COMPLEX, - CONV_THREAD, - CONV_OTHER, - CONV_INCLUDE, - CONV_INCLUDE_CUDA_MAIN_H, - CONV_TYPE, - CONV_LITERAL, - CONV_NUMERIC_LITERAL, - CONV_LAST + CONV_VERSION = 0, + CONV_INIT, + CONV_DEVICE, + CONV_MEM, + CONV_KERN, + CONV_COORD_FUNC, + CONV_MATH_FUNC, + CONV_DEVICE_FUNC, + CONV_SPECIAL_FUNC, + CONV_STREAM, + CONV_EVENT, + CONV_OCCUPANCY, + CONV_CONTEXT, + CONV_PEER, + CONV_MODULE, + CONV_CACHE, + CONV_EXEC, + CONV_EXTERNAL_RES, + CONV_GRAPH, + CONV_ERROR, + CONV_DEF, + CONV_TEX, + CONV_GL, + CONV_GRAPHICS, + CONV_SURFACE, + CONV_JIT, + CONV_D3D9, + CONV_D3D10, + CONV_D3D11, + CONV_VDPAU, + CONV_EGL, + CONV_COMPLEX, + CONV_THREAD, + CONV_OTHER, + CONV_INCLUDE, + CONV_INCLUDE_CUDA_MAIN_H, + CONV_TYPE, + CONV_LITERAL, + CONV_NUMERIC_LITERAL, + CONV_LAST }; constexpr int NUM_CONV_TYPES = (int) ConvTypes::CONV_LAST; enum ApiTypes { - API_DRIVER = 0, - API_RUNTIME, - API_BLAS, - API_RAND, - API_DNN, - API_FFT, - API_COMPLEX, - API_LAST + API_DRIVER = 0, + API_RUNTIME, + API_BLAS, + API_RAND, + API_DNN, + API_FFT, + API_COMPLEX, + API_LAST }; constexpr int NUM_API_TYPES = (int) ApiTypes::API_LAST; @@ -70,113 +70,81 @@ constexpr int NUM_API_TYPES = (int) ApiTypes::API_LAST; extern const char *counterNames[NUM_CONV_TYPES]; extern const char *apiNames[NUM_API_TYPES]; - struct hipCounter { - llvm::StringRef hipName; - ConvTypes type; - ApiTypes apiType; - bool unsupported; + llvm::StringRef hipName; + ConvTypes type; + ApiTypes apiType; + bool unsupported; }; - /** - * Tracks a set of named counters, as well as counters for each of the type enums defined above. - */ + * Tracks a set of named counters, as well as counters for each of the type enums defined above. + */ class StatCounter { private: - // Each thing we track is either "supported" or "unsupported"... - std::map counters; - - int apiCounters[NUM_API_TYPES] = {}; - int convTypeCounters[NUM_CONV_TYPES] = {}; + // Each thing we track is either "supported" or "unsupported"... + std::map counters; + int apiCounters[NUM_API_TYPES] = {}; + int convTypeCounters[NUM_CONV_TYPES] = {}; public: - void incrementCounter(const hipCounter& counter, std::string name); - - /** - * Add the counters from `other` onto the counters of this object. - */ - void add(const StatCounter& other); - - int getConvSum(); - - void print(std::ostream* csv, llvm::raw_ostream* printOut, std::string prefix); + void incrementCounter(const hipCounter& counter, std::string name); + // Add the counters from `other` onto the counters of this object. + void add(const StatCounter& other); + int getConvSum(); + void print(std::ostream* csv, llvm::raw_ostream* printOut, std::string prefix); }; /** - * Tracks the statistics for a single input file. - */ + * Tracks the statistics for a single input file. + */ class Statistics { - StatCounter supported; - StatCounter unsupported; - - std::string fileName; - - std::set touchedLines = {}; - int touchedBytes = 0; - - int totalLines = 0; - int totalBytes = 0; - - chr::steady_clock::time_point startTime; - chr::steady_clock::time_point completionTime; + StatCounter supported; + StatCounter unsupported; + std::string fileName; + std::set touchedLines = {}; + int touchedBytes = 0; + int totalLines = 0; + int totalBytes = 0; + chr::steady_clock::time_point startTime; + chr::steady_clock::time_point completionTime; public: - Statistics(std::string name); - - void incrementCounter(const hipCounter &counter, std::string name); - - /** - * Add the counters from `other` onto the counters of this object. - */ - void add(const Statistics &other); - - void lineTouched(int lineNumber); - void bytesChanged(int bytes); - - /** - * Set the completion timestamp to now. - */ - void markCompletion(); - - /////// Output functions /////// + Statistics(std::string name); + void incrementCounter(const hipCounter &counter, std::string name); + // Add the counters from `other` onto the counters of this object. + void add(const Statistics &other); + void lineTouched(int lineNumber); + void bytesChanged(int bytes); + // Set the completion timestamp to now. + void markCompletion(); public: - /** - * Pretty-print the statistics stored in this object. - * - * @param csv Pointer to an output stream for the CSV to write. If null, no CSV is written - * @param printOut Pointer to an output stream to print human-readable textual stats to. If null, no - * such stats are produced. - */ - void print(std::ostream* csv, llvm::raw_ostream* printOut, bool skipHeader = false); - - /// Print aggregated statistics for all registered counters. - static void printAggregate(std::ostream *csv, llvm::raw_ostream* printOut); - - /////// Static nonsense /////// - - // The Statistics for each input file. - static std::map stats; - - // The Statistics objects for the currently-being-processed input file. - static Statistics* currentStatistics; - - /** - * Aggregate statistics over all entries in `stats` and return the resulting Statistics object. - */ - static Statistics getAggregate(); - - /** - * Convenient global entry point for updating the "active" Statistics. Since we operate single-threadedly - * processing one file at a time, this allows us to simply expose the stats for the current file globally, - * simplifying things. - */ - static Statistics& current(); - - /** - * Set the active Statistics object to the named one, creating it if necessary, and write the completion - * timestamp into the currently active one. - */ - static void setActive(std::string name); + /** + * Pretty-print the statistics stored in this object. + * + * @param csv Pointer to an output stream for the CSV to write. If null, no CSV is written + * @param printOut Pointer to an output stream to print human-readable textual stats to. If null, no + * such stats are produced. + */ + void print(std::ostream* csv, llvm::raw_ostream* printOut, bool skipHeader = false); + // Print aggregated statistics for all registered counters. + static void printAggregate(std::ostream *csv, llvm::raw_ostream* printOut); + // The Statistics for each input file. + static std::map stats; + // The Statistics objects for the currently-being-processed input file. + static Statistics* currentStatistics; + // Aggregate statistics over all entries in `stats` and return the resulting Statistics object. + static Statistics getAggregate(); + /** + * Convenient global entry point for updating the "active" Statistics. Since we operate single-threadedly + * processing one file at a time, this allows us to simply expose the stats for the current file globally, + * simplifying things. + */ + static Statistics& current(); + /** + * Set the active Statistics object to the named one, creating it if necessary, and write the completion + * timestamp into the currently active one. + */ + static void setActive(std::string name); }; diff --git a/projects/hip/hipify-clang/src/StringUtils.cpp b/projects/hip/hipify-clang/src/StringUtils.cpp index ad55333bc8..6504d39010 100644 --- a/projects/hip/hipify-clang/src/StringUtils.cpp +++ b/projects/hip/hipify-clang/src/StringUtils.cpp @@ -1,17 +1,15 @@ #include "StringUtils.h" llvm::StringRef unquoteStr(llvm::StringRef s) { - if (s.size() > 1 && s.front() == '"' && s.back() == '"') { - return s.substr(1, s.size() - 2); - } - - return s; + if (s.size() > 1 && s.front() == '"' && s.back() == '"') { + return s.substr(1, s.size() - 2); + } + return s; } void removePrefixIfPresent(std::string &s, std::string prefix) { - if (s.find(prefix) != 0) { - return; - } - - s.erase(0, prefix.size()); + if (s.find(prefix) != 0) { + return; + } + s.erase(0, prefix.size()); } diff --git a/projects/hip/hipify-clang/src/StringUtils.h b/projects/hip/hipify-clang/src/StringUtils.h index 66a9be780f..c0be9f6227 100644 --- a/projects/hip/hipify-clang/src/StringUtils.h +++ b/projects/hip/hipify-clang/src/StringUtils.h @@ -4,11 +4,11 @@ #include "llvm/ADT/StringRef.h" /** - * Remove double-quotes from the start/end of a string, if present. - */ + * Remove double-quotes from the start/end of a string, if present. + */ llvm::StringRef unquoteStr(llvm::StringRef s); /** - * If `s` starts with `prefix`, remove it. Otherwise, does nothing. - */ + * If `s` starts with `prefix`, remove it. Otherwise, does nothing. + */ void removePrefixIfPresent(std::string &s, std::string prefix); diff --git a/projects/hip/hipify-clang/src/main.cpp b/projects/hip/hipify-clang/src/main.cpp index e420ab0681..4fdb8c87a7 100644 --- a/projects/hip/hipify-clang/src/main.cpp +++ b/projects/hip/hipify-clang/src/main.cpp @@ -1,5 +1,5 @@ /* -Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2015-2018 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 @@ -20,10 +20,11 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ /** - * @file Cuda2Hip.cpp - * - * This file is compiled and linked into clang based hipify tool. - */ + * @file Cuda2Hip.cpp + * + * This file is compiled and linked into clang based hipify tool. + */ + #include #include #include @@ -31,7 +32,6 @@ THE SOFTWARE. #include #include #include - #include "CUDA2HipMap.h" #include "LLVMCompat.h" #include "HipifyAction.h" @@ -42,7 +42,6 @@ THE SOFTWARE. namespace ct = clang::tooling; - namespace { void copyFile(const std::string& src, const std::string& dst) { @@ -55,7 +54,6 @@ void copyFile(const std::string& src, const std::string& dst) { int main(int argc, const char **argv) { llcompat::PrintStackTraceOnErrorSignal(); - ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::OneOrMore); std::vector fileSources = OptionsParser.getSourcePathList(); std::string dst = OutputFilename; @@ -63,7 +61,6 @@ int main(int argc, const char **argv) { llvm::errs() << "[HIPIFY] conflict: -o and multiple source files are specified.\n"; return 1; } - if (NoOutput) { if (Inplace) { llvm::errs() << "[HIPIFY] conflict: both -no-output and -inplace options are specified.\n"; @@ -74,13 +71,10 @@ int main(int argc, const char **argv) { return 1; } } - if (Examine) { NoOutput = PrintStats = true; } - int Result = 0; - // Arguments for the Statistics print routines. std::unique_ptr csv = nullptr; llvm::raw_ostream* statPrint = nullptr; @@ -90,7 +84,6 @@ int main(int argc, const char **argv) { if (PrintStats) { statPrint = &llvm::errs(); } - for (const auto & src : fileSources) { if (dst.empty()) { if (Inplace) { @@ -102,55 +95,42 @@ int main(int argc, const char **argv) { llvm::errs() << "[HIPIFY] conflict: both -o and -inplace options are specified.\n"; return 1; } - std::string tmpFile = src + ".hipify-tmp"; - // Create a copy of the file to work on. When we're done, we'll move this onto the // output (which may mean overwriting the input, if we're in-place). // Should we fail for some reason, we'll just leak this file and not corrupt the input. copyFile(src, tmpFile); - // Initialise the statistics counters for this file. Statistics::setActive(src); - // RefactoringTool operates on the file in-place. Giving it the output path is no good, // because that'll break relative includes, and we don't want to overwrite the input file. // So what we do is operate on a copy, which we then move to the output. ct::RefactoringTool Tool(OptionsParser.getCompilations(), tmpFile); ct::Replacements& replacementsToUse = llcompat::getReplacements(Tool, tmpFile); - ReplacementsFrontendActionFactory actionFactory(&replacementsToUse); - Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("--cuda-host-only", ct::ArgumentInsertPosition::BEGIN)); - // Ensure at least c++11 is used. Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-std=c++11", ct::ArgumentInsertPosition::BEGIN)); #if defined(HIPIFY_CLANG_RES) Tool.appendArgumentsAdjuster(ct::getInsertArgumentAdjuster("-resource-dir=" HIPIFY_CLANG_RES)); #endif Tool.appendArgumentsAdjuster(ct::getClangSyntaxOnlyAdjuster()); - // Hipify _all_ the things! if (Tool.runAndSave(&actionFactory)) { LLVM_DEBUG(llvm::dbgs() << "Skipped some replacements.\n"); } - // Either move the tmpfile to the output, or remove it. if (!NoOutput) { rename(tmpFile.c_str(), dst.c_str()); } else { remove(tmpFile.c_str()); } - Statistics::current().markCompletion(); Statistics::current().print(csv.get(), statPrint); - dst.clear(); } - if (fileSources.size() > 1) { Statistics::printAggregate(csv.get(), statPrint); } - return Result; } From 68ba937cf72eb2f1fa320dd6e0e4121c85bd813c Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 17 Oct 2018 11:03:43 +0530 Subject: [PATCH 8/9] Update Jenkinsfile Disable rocm-1.9x testing in CI due to incompatible changes with HCC from ROCm 1.9.x [ROCm/hip commit: f5356fa9ecee3e9c2e667db160e9b8008c6db4d8] --- projects/hip/Jenkinsfile | 20 +++++++++----------- 1 file changed, 9 insertions(+), 11 deletions(-) diff --git a/projects/hip/Jenkinsfile b/projects/hip/Jenkinsfile index ab6dd6d67f..24f2c4ae44 100644 --- a/projects/hip/Jenkinsfile +++ b/projects/hip/Jenkinsfile @@ -297,7 +297,7 @@ String build_config = 'Release' String job_name = env.JOB_NAME.toLowerCase( ) // The following launches 3 builds in parallel: hcc-ctu, hcc-1.6 and cuda -parallel rocm_1_9: +parallel /*rocm_1_9: { node('hip-rocm') { @@ -331,17 +331,15 @@ parallel rocm_1_9: docker_clean_images( 'hip', docker_build_image_name( ) ) // After a successful build, upload a docker image of the results - /* - String hip_image_name = docker_upload_artifactory( hcc_ver, job_name, from_image, source_hip_rel, build_hip_rel ) - if( params.push_image_to_docker_hub ) - { - docker_upload_dockerhub( job_name, hip_image_name, 'rocm' ) - docker_clean_images( 'rocm', hip_image_name ) - } - docker_clean_images( job_name, hip_image_name ) - */ + //String hip_image_name = docker_upload_artifactory( hcc_ver, job_name, from_image, source_hip_rel, build_hip_rel ) + //if( params.push_image_to_docker_hub ) + //{ + // docker_upload_dockerhub( job_name, hip_image_name, 'rocm' ) + // docker_clean_images( 'rocm', hip_image_name ) + //} + //docker_clean_images( job_name, hip_image_name ) } -}, +},*/ rocm_head: { node('hip-rocm') From c6905e9e6b6f040aac6fbba78a5f34a1b1be1c59 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 17 Oct 2018 11:13:37 +0530 Subject: [PATCH 9/9] Fix typos in Jenkinsfile [ROCm/hip commit: c9e4ed1149a96a59657fc52e94acb915d4c05822] --- projects/hip/Jenkinsfile | 47 ++-------------------------------------- 1 file changed, 2 insertions(+), 45 deletions(-) diff --git a/projects/hip/Jenkinsfile b/projects/hip/Jenkinsfile index 24f2c4ae44..540e7e6689 100644 --- a/projects/hip/Jenkinsfile +++ b/projects/hip/Jenkinsfile @@ -296,51 +296,8 @@ def docker_upload_dockerhub( String local_org, String image_name, String remote_ String build_config = 'Release' String job_name = env.JOB_NAME.toLowerCase( ) -// The following launches 3 builds in parallel: hcc-ctu, hcc-1.6 and cuda -parallel /*rocm_1_9: -{ - node('hip-rocm') - { - String hcc_ver = 'rocm-1.9.x' - String from_image = 'ci_test_nodes/rocm-1.9.x/ubuntu-16.04:latest' - String inside_args = '--device=/dev/kfd --device=/dev/dri --group-add=video' - - // Checkout source code, dependencies and version files - String source_hip_rel = checkout_and_version( hcc_ver ) - - // Create/reuse a docker image that represents the hip build environment - def hip_build_image = docker_build_image( hcc_ver, 'hip', '', source_hip_rel, from_image ) - - // Print system information for the log - hip_build_image.inside( inside_args ) - { - sh """#!/usr/bin/env bash - set -x - /opt/rocm/bin/rocm_agent_enumerator -t ALL - /opt/rocm/bin/hcc --version - """ - } - - // Conctruct a binary directory path based on build config - String build_hip_rel = build_directory_rel( build_config ); - - // Build hip inside of the build environment - docker_build_inside_image( hip_build_image, inside_args, hcc_ver, '', build_config, source_hip_rel, build_hip_rel ) - - // Clean docker build image - docker_clean_images( 'hip', docker_build_image_name( ) ) - - // After a successful build, upload a docker image of the results - //String hip_image_name = docker_upload_artifactory( hcc_ver, job_name, from_image, source_hip_rel, build_hip_rel ) - //if( params.push_image_to_docker_hub ) - //{ - // docker_upload_dockerhub( job_name, hip_image_name, 'rocm' ) - // docker_clean_images( 'rocm', hip_image_name ) - //} - //docker_clean_images( job_name, hip_image_name ) - } -},*/ -rocm_head: +// The following launches 2 builds in parallel: rocm-head and cuda-9.x +parallel rocm_head: { node('hip-rocm') {