ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-18 20:09:25 +00:00
Code Issues Packages Projects Releases Wiki Activity

Change all device operations to use add_instance_library (#338)

Browse Source 
* Change all device operations to use add_instance_library to avoid duplicated cmake configuration.

* update DeviceMem

Co-authored-by: Chao Liu <chao.liu2@amd.com>

[ROCm/composable_kernel commit: fb1cbf025b]
This commit is contained in:
cloudhan
2022-08-14 01:17:58 +08:00
committed by GitHub
parent d15689b96a
commit f1a2efceb5
31 changed files with 190 additions and 358 deletions
Download Patch File Download Diff File Expand all files Collapse all files

31
library/include/ck/library/utility/device_memory.hpp
View File

@@ -18,23 +18,26 @@ struct DeviceMem
{
DeviceMem() = delete;
DeviceMem(std::size_t mem_size);
void* GetDeviceBuffer();
std::size_t GetBufferSize();
void ToDevice(const void* p);
void FromDevice(void* p);
void SetZero();
void* GetDeviceBuffer() const;
std::size_t GetBufferSize() const;
void ToDevice(const void* p) const;
void FromDevice(void* p) const;
void SetZero() const;
template <typename T>
void SetValue(T x)
{
if(mMemSize % sizeof(T) != 0)
{
throw std::runtime_error("wrong! not entire DeviceMem will be set");
}
set_buffer_value<T><<<1, 1024>>>(static_cast<T*>(mpDeviceBuf), x, mMemSize / sizeof(T));
}
void SetValue(T x) const;
~DeviceMem();
void* mpDeviceBuf;
std::size_t mMemSize;
};
template <typename T>
void DeviceMem::SetValue(T x) const
{
if(mMemSize % sizeof(T) != 0)
{
throw std::runtime_error("wrong! not entire DeviceMem will be set");
}
set_buffer_value<T><<<1, 1024>>>(static_cast<T*>(mpDeviceBuf), x, mMemSize / sizeof(T));
}