From 20847d2cb421a0968d2abbc4baddfd5e51a7c365 Mon Sep 17 00:00:00 2001 From: James Osborn Date: Sun, 17 Sep 2023 20:03:08 -0500 Subject: [PATCH] add missing file remove some unused arguments --- include/kernels/restrictor.cuh | 2 +- include/targets/generic/helpers.h | 2 +- include/targets/generic/special_ops.h | 18 ++--- include/targets/hip/shared_memory_helper.h | 93 ++++++++++++++++++++++ include/targets/hip/special_ops_target.h | 6 +- 5 files changed, 107 insertions(+), 14 deletions(-) create mode 100644 include/targets/hip/shared_memory_helper.h diff --git a/include/kernels/restrictor.cuh b/include/kernels/restrictor.cuh index b4aaf60c9d..59810f174c 100644 --- a/include/kernels/restrictor.cuh +++ b/include/kernels/restrictor.cuh @@ -119,7 +119,7 @@ namespace quda { static constexpr const char *filename() { return KERNEL_FILE; } template - __device__ __host__ inline void operator()(dim3 block, dim3 thread, bool active = true) + __device__ __host__ inline void operator()(dim3 block, dim3 thread, bool /*active*/ = true) { int x_fine_offset = thread.x; const int x_coarse = block.x; diff --git a/include/targets/generic/helpers.h b/include/targets/generic/helpers.h index 9f59ee5243..eeb6a65533 100644 --- a/include/targets/generic/helpers.h +++ b/include/targets/generic/helpers.h @@ -45,7 +45,7 @@ namespace quda template struct DimsStatic { - static constexpr dim3 dims(dim3 block) { + static constexpr dim3 dims(dim3) { return dim3(x,y,z); } }; diff --git a/include/targets/generic/special_ops.h b/include/targets/generic/special_ops.h index 563f48af61..83498e57c5 100644 --- a/include/targets/generic/special_ops.h +++ b/include/targets/generic/special_ops.h @@ -6,25 +6,25 @@ namespace quda { // dimensions functors for SharedMemoryCache struct opDimsBlock { - template static constexpr dim3 dims(dim3 b, const Arg &...arg) { return b; } + template static constexpr dim3 dims(dim3 b, const Arg &...) { return b; } }; template struct opDimsStatic { - template static constexpr dim3 dims(dim3 b, const Arg &...arg) { return dim3(bx,by,bz); } + template static constexpr dim3 dims(dim3, const Arg &...) { return dim3(bx,by,bz); } }; // size functors for determining shared memory size struct opSizeBlock { - template static constexpr unsigned int size(dim3 b, const Arg &...arg) { + template static constexpr unsigned int size(dim3 b, const Arg &...) { return b.x * b.y * b.z * sizeof(T); } }; struct opSizeBlockDivWarp { - template static constexpr unsigned int size(dim3 b, const Arg &...arg) { + template static constexpr unsigned int size(dim3 b, const Arg &...) { return ((b.x * b.y * b.z + device::warp_size() - 1)/device::warp_size()) * sizeof(T); } }; template struct opSizeStatic { - template static constexpr unsigned int size(dim3 b, const Arg &...arg) { + template static constexpr unsigned int size(dim3 b, const Arg &...) { return S * sizeof(T); } }; @@ -54,7 +54,7 @@ namespace quda { static constexpr bool hasSpecialOp> = ( std::is_same_v || ... ); //template void checkSpecialOps() { static_assert(hasSpecialOp); } - template void checkSpecialOps(const Ops &ops) { + template void checkSpecialOps(const Ops &) { static_assert(hasSpecialOp); } @@ -248,7 +248,7 @@ namespace quda { }; template <> struct sharedMemSizeS { template - static constexpr unsigned int size(dim3, Arg &...arg) { + static constexpr unsigned int size(dim3, Arg &...) { return 0; } }; @@ -266,7 +266,7 @@ namespace quda { // sharedMemOffset template struct sharedMemOffset { template - inline int operator()(dim3 block, Arg &...arg) { return 0; } + inline int operator()(dim3 block, Arg &...) { return 0; } }; template struct sharedMemOffset,n> { template @@ -274,7 +274,7 @@ namespace quda { }; template struct sharedMemOffset,0> { template - inline int operator()(dim3 block, Arg &...arg) { return 0; } + inline int operator()(dim3 block, Arg &...) { return 0; } }; template struct sharedMemOffset,n> { template diff --git a/include/targets/hip/shared_memory_helper.h b/include/targets/hip/shared_memory_helper.h new file mode 100644 index 0000000000..c29d14eb11 --- /dev/null +++ b/include/targets/hip/shared_memory_helper.h @@ -0,0 +1,93 @@ +#pragma once + +#include +#include + +/** + @file shared_memory_helper.h + + Target specific helper for allocating and accessing shared memory. + */ + +namespace quda +{ + + /** + @brief Class which is used to allocate and access shared memory. + The shared memory is treated as an array of type T, with the + number of elements given by a call to the static member + S::size(target::block_dim()). The byte offset from the beginning + of the total shared memory block is given by the static member + O::shared_mem_size(target::block_dim()), or 0 if O is void. + */ + template class SharedMemory + { + public: + using value_type = T; + + private: + T *data; + + /** + @brief This is a dummy instantiation for the host compiler + */ + template struct cache_dynamic { + T *operator()(unsigned int) + { + static T *cache_; + return cache_; + } + }; + + /** + @brief This is the handle to the dynamic shared memory + @return Shared memory pointer + */ + template struct cache_dynamic { + __device__ inline T *operator()(unsigned int offset) + { + extern __shared__ int cache_[]; + return reinterpret_cast(reinterpret_cast(cache_) + offset); + } + }; + + __device__ __host__ inline T *cache(unsigned int offset) const { return target::dispatch(offset); } + + public: + /** + @brief Byte offset for this shared memory object. + */ + static constexpr unsigned int get_offset(dim3 block) + { + unsigned int o = 0; + if constexpr (!std::is_same_v) { o = O::shared_mem_size(block); } + return o; + } + + /** + @brief Shared memory size in bytes. + */ + static constexpr unsigned int shared_mem_size(dim3 block) { return get_offset(block) + S::size(block) * sizeof(T); } + + /** + @brief Constructor for SharedMemory object. + */ + HostDevice constexpr SharedMemory() : data(cache(get_offset(target::block_dim()))) { } + + template + HostDevice constexpr SharedMemory(const SpecialOps &ops) : data(cache(get_offset(target::block_dim()))) { } + + /** + @brief Return this SharedMemory object. + */ + HostDevice constexpr auto sharedMem() const { return *this; } + + /** + @brief Subscripting operator returning a reference to element. + @param[in] i The index to use. + @return Reference to value stored at that index. + */ + HostDevice T &operator[](const int i) const { return data[i]; } + }; + +} // namespace quda diff --git a/include/targets/hip/special_ops_target.h b/include/targets/hip/special_ops_target.h index dbc4703d83..0975cbfdbb 100644 --- a/include/targets/hip/special_ops_target.h +++ b/include/targets/hip/special_ops_target.h @@ -6,7 +6,7 @@ namespace quda { // SpecialOps template struct SpecialOps : SpecialOps_Base { - template constexpr void setSpecialOps(const SpecialOps &ops) { + template constexpr void setSpecialOps(const SpecialOps &) { static_assert(std::is_same_v,SpecialOps>); } }; @@ -15,7 +15,7 @@ namespace quda { struct op_blockSync : op_BaseT { //using dependencies = depFullBlock; template - static constexpr unsigned int shared_mem_size(dim3 block, Arg &...arg) { return 0; } + static constexpr unsigned int shared_mem_size(dim3 block, Arg &...) { return 0; } }; template @@ -23,7 +23,7 @@ namespace quda { //using dependencies = depNone; //using dependencies = depFullBlock; template - static constexpr unsigned int shared_mem_size(dim3, Arg &...arg) { return 0; } + static constexpr unsigned int shared_mem_size(dim3, Arg &...) { return 0; } }; }