mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-20 04:49:54 +00:00
8cba08d07a
* add gridwise gemm v4r1
* rename
* adding gemm+reduce
* adding gemm+reduce
* adding gemm+reduce
* adding gemm+reduce
* use sfc in shuffling
* remove hardcode
* remove hardcode
* refactor
* fix build
* adding gemm+reduce
* adding gemm+reduce
* adding gemm+reduce
* adding gemm+reduce
* adding gemm+reduce
* format
* clean
* adding gemm+reduce
* adding profiler for gemm+reduce
* adding gemm+reduce profiler
* fix build
* clean up
* gemm+reduce
* fix build
* update DeviceGemm_Xdl_CShuffle; update enum to enum class
* clean up
* add test for gemm+reduce
* clean up
* refactor
* fix build
* fix build
[ROCm/composable_kernel commit: f95267f166]
45 lines
1.1 KiB
C++
45 lines
1.1 KiB
C++
#ifndef CK_AMD_ADDRESS_SPACE_HPP
|
|
#define CK_AMD_ADDRESS_SPACE_HPP
|
|
|
|
#include "config.hpp"
|
|
#include "c_style_pointer_cast.hpp"
|
|
|
|
// Address Space for AMDGCN
|
|
// https://llvm.org/docs/AMDGPUUsage.html#address-space
|
|
|
|
namespace ck {
|
|
|
|
enum struct AddressSpaceEnum_t
|
|
{
|
|
Generic,
|
|
Global,
|
|
Lds,
|
|
Sgpr,
|
|
Vgpr,
|
|
};
|
|
|
|
template <typename T>
|
|
__device__ T* cast_pointer_to_generic_address_space(T CONSTANT* p)
|
|
{
|
|
// cast a pointer in "Constant" address space (4) to "Generic" address space (0)
|
|
// only c-style pointer cast seems be able to be compiled
|
|
#pragma clang diagnostic push
|
|
#pragma clang diagnostic ignored "-Wold-style-cast"
|
|
return (T*)p; // NOLINT(old-style-cast)
|
|
#pragma clang diagnostic pop
|
|
}
|
|
|
|
template <typename T>
|
|
__host__ __device__ T CONSTANT* cast_pointer_to_constant_address_space(T* p)
|
|
{
|
|
// cast a pointer in "Generic" address space (0) to "Constant" address space (4)
|
|
// only c-style pointer cast seems be able to be compiled
|
|
#pragma clang diagnostic push
|
|
#pragma clang diagnostic ignored "-Wold-style-cast"
|
|
return (T CONSTANT*)p; // NOLINT(old-style-cast)
|
|
#pragma clang diagnostic pop
|
|
}
|
|
|
|
} // namespace ck
|
|
#endif
|