Merge pull request llvm#301 from AMD-Lightning-Internal/amd/dev/bsumner_amdeng/cherrypick_id_mask_fix

Amd/dev/bsumner amdeng/cherrypick id mask fix

Author: ronlieb

amd/device-libs/ockl/src/hsaqs.cl

@@ -72,7 +72,11 @@ update_mbox(const __global amd_signal_t *sig)
     if (mb) {
         uint id = sig->event_id;
         AS(mb, id, memory_order_release, memory_scope_all_svm_devices);
-        __builtin_amdgcn_s_sendmsg(1 | (0 << 4), __builtin_amdgcn_readfirstlane(id) & 0xff);
+        uint mid = id &
+                    (__oclc_ISA_version < 9000 ? 0xff :
+                     (__oclc_ISA_version < 10000 ? 0xffffff :
+                      (__oclc_ISA_version < 11000 ? 0x7fffff : 0xffffff)));
+        __builtin_amdgcn_s_sendmsg(1 | (0 << 4), __builtin_amdgcn_readfirstlane(mid));
     }
 }
 

offload/hostexec/src/hostexec_invoke.cpp

@@ -1,5 +1,6 @@
 
 #include <stdint.h>
+#include "Platform.h"
 
 #define GLOB_ATTR __attribute__((address_space(1)))
 #define __static_inl static __attribute__((flatten, always_inline))