#include "Interface.h"
#include "Synchronization.h"
#include "Types.h"
+#include "Utils.h"
using namespace _OMP;
void SharedMemorySmartStackTy::pop(void *Ptr, uint32_t Bytes) {
uint64_t AlignedBytes = utils::align_up(Bytes, Alignment);
- if (Ptr >= &Data[0] && Ptr < &Data[state::SharedScratchpadSize]) {
+ if (utils::isSharedMemPtr(Ptr)) {
int TId = mapping::getThreadIdInBlock();
Usage[TId] -= AlignedBytes;
return;
namespace impl {
+bool isSharedMemPtr(const void *Ptr) { return false; }
void Unpack(uint64_t Val, uint32_t *LowBits, uint32_t *HighBits);
uint64_t Pack(uint32_t LowBits, uint32_t HighBits);
}
#pragma omp end declare variant
+///}
/// NVPTX Implementation
///
}
#pragma omp end declare variant
+///}
int32_t shuffle(uint64_t Mask, int32_t Var, int32_t SrcLane);
int32_t shuffleDown(uint64_t Mask, int32_t Var, uint32_t LaneDelta,
return __builtin_amdgcn_ds_bpermute(Index << 2, Var);
}
+bool isSharedMemPtr(const void * Ptr) {
+ return __builtin_amdgcn_is_shared((const __attribute__((address_space(0))) void *)Ptr);
+}
#pragma omp end declare variant
///}
return __nvvm_shfl_sync_down_i32(Mask, Var, Delta, T);
}
+bool isSharedMemPtr(const void *Ptr) { return __nvvm_isspacep_shared(Ptr); }
+
#pragma omp end declare variant
+///}
} // namespace impl
uint64_t utils::pack(uint32_t LowBits, uint32_t HighBits) {
return impl::shuffleDown(Mask, Var, Delta, Width);
}
+bool utils::isSharedMemPtr(void *Ptr) { return impl::isSharedMemPtr(Ptr); }
+
extern "C" {
int32_t __kmpc_shuffle_int32(int32_t Val, int16_t Delta, int16_t SrcLane) {
FunctionTracingRAII();