Add hipclang amdgcn functions by aaronenyeshi · Pull Request #515 · ROCm/hip

aaronenyeshi · 2018-06-12T22:08:47Z

These are moving from hipclang in device library to hip headers. These are required for the functionality of HIPclang project.

aaronenyeshi · 2018-06-12T22:10:16Z

So far I have added assertfail, ballot, clock, lanemask, memory, named sync, smid, sync, and trap. These previously existed as .cl and .ll files in https://github.com/RadeonOpenCompute/ROCm-Device-Libs/tree/9b46f23b1795b56600bb39ff2594ce129ab0222e/hip/src

yxsamliu · 2018-06-13T11:16:04Z

+
+// hip.amdgcn.bc - sync threads
+// extern "C" __device__ __attribute__((noduplicate)) void __syncthreads();
+#define CLK_LOCAL_MEM_FENCE    0x01


please prefix with __, since we don't want to pollute users' name space

yxsamliu · 2018-06-13T11:16:43Z

+// hip.amdgcn.bc - sync threads
+// extern "C" __device__ __attribute__((noduplicate)) void __syncthreads();
+#define CLK_LOCAL_MEM_FENCE    0x01
+#define local __attribute__((address_space(3)))


prefix local with __

yxsamliu · 2018-06-13T11:17:06Z

+#define CLK_LOCAL_MEM_FENCE    0x01
+#define local __attribute__((address_space(3)))
+
+typedef unsigned cl_mem_fence_flags;


prefix with __

yxsamliu · 2018-06-13T11:17:15Z

+
+typedef unsigned cl_mem_fence_flags;
+
+typedef enum memory_scope {


prefix with __

yxsamliu · 2018-06-13T11:19:06Z

+} memory_scope;
+
+// enum values aligned with what clang uses in EmitAtomicExpr()
+typedef enum memory_order


prefix with __

yxsamliu · 2018-06-13T11:22:55Z

+} memory_order;
+
+extern "C" __device__ __attribute__((overloadable))
+void atomic_work_item_fence(cl_mem_fence_flags, memory_order, memory_scope);


we cannot introduce this function name without prefix __. we need either implement hc_work_group_barrier with functions the name of which is prefixed with __, or add alias prefixed with __ to atomic_work_item_fence in the device library.

yxsamliu · 2018-06-13T16:51:55Z


+__device__
+inline
+void* __get_dynamicgroupbaseptr() { return get_dynamic_group_segment_base_pointer(); }


can we remove get_dynamic_group_segment_base_pointer and implement __get_dynamicgroupbaseptr directly? I think the disabled code in __amdgcn_get_dynamicgroupbaseptr is related to this.

yxsamliu · 2018-06-13T16:52:52Z

+
+__device__
+inline
+static void hc_work_group_barrier(__cl_mem_fence_flags flags, __memory_scope scope)


can we rename it to __work_group_barrier?

yxsamliu · 2018-06-13T16:54:01Z

+
+__device__
+inline
+static void hc_barrier(int n)


can we rename it as __barrier?

yxsamliu · 2018-06-13T17:08:56Z

I think it is cleaner to create two new header files:

llvm_intrinsics.h - containing declarations for wrapper functions for llvm intrinsics like llvm.amdgcn.s.barrier

device_library_decls.h - containing declarations for types and functions in device library

Other changes in this patch can be moved to device_functions.h

aaronenyeshi · 2018-06-13T20:42:37Z

I'm working on separating this into 3 files, llvm_intrinsics.h, device_library_decls.h, and device_functions.h. Thanks @yxsamliu

yxsamliu

LGTM. Thanks!

These are moving from hipclang in device library to hip headers. These are required for the functionality of HIPclang project.

Move all Integer Intrinsics, device_functions.cpp definitions and HIP specific device functions into HIP headers. Implement the device functions using llvm_intrinsics and device-libs functions instead of calling hc::__* functions. Remove device_functions.cpp since everything is now defined in header.

Move all __hip_hc_ir_* functions from hip_hc.ll into HIP header as inline asm. Remove hip_hc.ll and build dependencies from HIP.

Fix surface test on HIP clang path.

aaronenyeshi · 2018-07-05T21:38:33Z

@mangupta I've added a few more changes to allow hip-clang path to pass tests: hipSurfaceObj2D, hip_test_ldg, hip_threadfence_system, and hipNullStream. Thanks!

yxsamliu · 2018-07-08T03:14:11Z

+    return ballot;
+}
+
+#ifdef __HIP_DEVICE_COMPILE__


we still need the device functions to be declared for host compilation, therefore the code should be like this

#ifdef __HIP_DEVICE_COMPILE__ __device__ inline void* __get_dynamicgroupbaseptr() { // Get group segment base pointer. return (char*)__local_to_generic(__to_local(__llvm_amdgcn_groupstaticsize())); } #else __device__ void* __get_dynamicgroupbaseptr(); #endif // __HIP_DEVICE_COMPILE__

Hi Sam, I've cherry-picked your fixes into this PR related to your review. Thanks!

aaronenyeshi · 2018-07-10T14:44:14Z

@yxsamliu @scchan @mangupta I believe this is ready to land. Please help review. It passes all unit tests on HCC/HIP, also it passes 97/106 unit tests on HIP-clang on HCC rt. Thank you!

yxsamliu · 2018-07-10T15:15:20Z

Pleas hold on. We need Brian to review it. Thanks.

aaronenyeshi · 2018-07-10T15:48:30Z

@b-sumner - I will make a new commit to remove the old usage of __activelanemake_v4_b64_b1.

Replace implementation of __any and __all functions using OCKL functions and replaced __ballot implementation to use llvm intrinsic llvm.amdgcn.icmp.i32 instead of calls to __activelanemask_v4_b64_b1 which is not convergent.

b-sumner · 2018-07-10T17:58:33Z

+
+extern "C" __device__ int32_t __ockl_activelane_u32(void);
+
+extern "C" __device__ uint __ockl_mul24_u32(uint, uint);


Please copy the function attributes these functions are given in ocml.h.

b-sumner · 2018-07-10T18:02:10Z

-//__device__ void  __threadfence_system(void) __attribute__((deprecated("Provided with workaround
-//configuration, see hip_kernel_language.md for details")));
-__device__ void __threadfence_system(void);
+__device__ inline static int min(int arg1, int arg2) {


These can't work. An f32 can not exactly represent every i32. Just use arg1 < arg2 ? arg1 : arg2. The compiler will do the rest.

b-sumner · 2018-07-10T18:04:30Z

+
+#include "hip/hcc_detail/host_defines.h"
+
+__device__ ulong __llvm_amdgcn_icmp_i32(uint x, uint y, uint z) __asm("llvm.amdgcn.icmp.i32");


This should have the convergent attribute if it is really needed.

b-sumner · 2018-07-10T18:17:18Z

+__device__ static inline int __float2int_ru(float x) { return (int)__ocml_ceil_f32(x); }
+__device__ static inline int __float2int_rz(float x) { return (int)__ocml_trunc_f32(x); }
+
+__device__ static inline long long int __float2ll_rd(float x) { return (long long int)x; }


These are not correct. There should at least be a comment indicating that. Are correct implementations needed?

I copied these conversion functions from old HIP implementation. How can I fix it?

I think the correct versions cause isel failures in backend. Probably just add ToDo.

b-sumner · 2018-07-10T18:20:05Z

+
+__device__ static inline double __int2double_rn(int x) { return (double)x; }
+
+__device__ static inline float __int2float_rd(int x) { return (float)x; }


These are not correct. Do they need to be?

b-sumner · 2018-07-10T18:26:33Z

I only noted a few of the conversion function problems. There are several others.

aaronenyeshi · 2018-07-10T18:59:00Z

Thanks @b-sumner, I've added the missing attributes, changed llvm.amdgcn.icmp.i32 asm and changed min/max as you suggested. For the conversion functions, I've left a TODO comment for the future.

changed __to_local

Hi Brian, I've revised to your reviews. Let me know if it is resolved. Also we've discussed that we can keep the conversion functions as it is for now with a TODO.

b-sumner · 2018-07-12T20:12:46Z

Looks good to me. I will work on the conversion functions after this is checked in.

aaronenyeshi · 2018-07-16T15:08:54Z

@mangupta ping - we need this merged for internal HIP testing and investigation.

aaronenyeshi requested a review from yxsamliu June 12, 2018 22:08

yxsamliu reviewed Jun 13, 2018

View reviewed changes

aaronenyeshi closed this Jun 13, 2018

aaronenyeshi reopened this Jun 13, 2018

aaronenyeshi force-pushed the hipclang-add-amdgcn-funcs branch 3 times, most recently from f74ce78 to 219e717 Compare June 18, 2018 22:46

yxsamliu previously approved these changes Jun 19, 2018

View reviewed changes

aaronenyeshi requested a review from mangupta June 19, 2018 17:15

aaronenyeshi dismissed yxsamliu’s stale review via 33417b7 June 20, 2018 17:33

aaronenyeshi added 9 commits June 20, 2018 17:38

Add hipclang amdgcn functions

c453b42

These are moving from hipclang in device library to hip headers. These are required for the functionality of HIPclang project.

Add __llvm_fence funcs and __ prefixes

cfa8fc1

Add prefix __ to memory scope and order

871cfc6

Add get_dynamicgroupbaseptr def and remove hc_

27f600b

Split __llvm and device lib funcs into new headers

cfe3748

Move hipclang funcs into corresponding headers

fe4e6c5

Implement hip_hc.ll into HIP headers

2142eb4

Move all __hip_hc_ir_* functions from hip_hc.ll into HIP header as inline asm. Remove hip_hc.ll and build dependencies from HIP.

Implement __ballot, __any, __all into HIP headers

6dc16bb

Enable surface functions on HIP clang path

22b60af

Fix surface test on HIP clang path.

yxsamliu reviewed Jul 8, 2018

View reviewed changes

yxsamliu added 2 commits July 9, 2018 17:56

Add declare of __get_dynamicgroupbaseptr for host compilation

e0390dd

Add workaround __local_to_generic

5ec7973

mangupta mentioned this pull request Jul 10, 2018

Adding missing __shfl overloads. #557

Closed

Remove activelanemask asm using ockl and llvm instrinsics

d1323f4

Replace implementation of __any and __all functions using OCKL functions and replaced __ballot implementation to use llvm intrinsic llvm.amdgcn.icmp.i32 instead of calls to __activelanemask_v4_b64_b1 which is not convergent.

b-sumner requested changes Jul 10, 2018

View reviewed changes

b-sumner previously requested changes Jul 10, 2018

View reviewed changes

aaronenyeshi added 2 commits July 10, 2018 18:27

Add func attributes to match ocml and ockl

1f87173

Fix min/max, icmp asm and add comment for conversion functions

4a7ad93

mangupta added the pr:ready_for_ci PR ready to be tested label Jul 11, 2018

mangupta approved these changes Jul 17, 2018

View reviewed changes

mangupta merged commit fbbe259 into master Jul 17, 2018

aaronenyeshi deleted the hipclang-add-amdgcn-funcs branch July 17, 2018 18:28

This was referenced Jul 18, 2018

No support for __shfl_xor or __shfl_down for type double. #434

Closed

Request to expose __to_local() to HIP kernels #430

Closed

shfl operations for fp64 type. #442

Closed

jszuppe mentioned this pull request Aug 22, 2018

Initial changes for HIP-Clang Compiler ROCm/rocPRIM#23

Closed


		typedef unsigned cl_mem_fence_flags;

		typedef enum memory_scope {


		extern "C" __device__ int32_t __ockl_activelane_u32(void);

		extern "C" __device__ uint __ockl_mul24_u32(uint, uint);


		#include "hip/hcc_detail/host_defines.h"

		__device__ ulong __llvm_amdgcn_icmp_i32(uint x, uint y, uint z) __asm("llvm.amdgcn.icmp.i32");


		__device__ static inline double __int2double_rn(int x) { return (double)x; }

		__device__ static inline float __int2float_rd(int x) { return (float)x; }

Conversation

aaronenyeshi commented Jun 12, 2018

Uh oh!

aaronenyeshi commented Jun 12, 2018

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

yxsamliu commented Jun 13, 2018

Uh oh!

aaronenyeshi commented Jun 13, 2018

Uh oh!

yxsamliu left a comment

Choose a reason for hiding this comment

Uh oh!

aaronenyeshi commented Jul 5, 2018

Uh oh!

yxsamliu Jul 8, 2018 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

aaronenyeshi commented Jul 10, 2018

Uh oh!

yxsamliu commented Jul 10, 2018

Uh oh!

aaronenyeshi commented Jul 10, 2018

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

b-sumner commented Jul 10, 2018

Uh oh!

aaronenyeshi commented Jul 10, 2018

Uh oh!

b-sumner commented Jul 12, 2018

Uh oh!

aaronenyeshi commented Jul 16, 2018

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

yxsamliu Jul 8, 2018 •

edited

Loading