ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-04-20 14:59:17 +00:00
Code Issues Packages Projects Releases Wiki Activity

Compile CK for all targets (#188)

Browse Source 
* compile ck for all targets

* update the target criteria

* change the target condition

* fixed some typos

* fixed missed file

* revert changes in README

* revert device_conv3d_fwd_xdl_...

* update device_conv3d_fwd_xdl_...

* update device_batched_gemm_reduce...

* test the unused arguments fix

* test the warning suppression

* try suppress warnings in device_batched_gemm_reduce_xdl...

* fix the last warnings

* replace UNUSED with std::ignore

* fix a typo

* replaced std::ignore with ignore

* add igonre header to common_header

* refactor atomicAdd

Co-authored-by: Chao Liu <chao.liu2@amd.com>
This commit is contained in:
Illia Silin
2022-04-15 12:17:28 -07:00
committed by GitHub
parent ac0d806650
commit 4221505d3e
17 changed files with 232 additions and 77 deletions
Download Patch File Download Diff File Expand all files Collapse all files

13
include/ck/utility/dynamic_buffer.hpp
View File

@@ -1,11 +1,16 @@
#pragma once
#include "amd_buffer_addressing.hpp"
#include "c_style_pointer_cast.hpp"
#include "config.hpp"
#include "enable_if.hpp"
#include "c_style_pointer_cast.hpp"
#include "amd_buffer_addressing.hpp"
#include "generic_memory_space_atomic_add.hpp"
namespace ck {
// T may be scalar or vector
// X may be scalar or vector
// T and X have same scalar type
// X contains multiple T
template <AddressSpaceEnum BufferAddressSpace,
typename T,
typename ElementSpaceSize,
@@ -316,9 +321,7 @@ struct DynamicBuffer
{
if(is_valid_element)
{
// FIXME: atomicAdd is defined by HIP, need to avoid implicit type casting when
// calling it
atomicAdd(c_style_pointer_cast<X*>(&p_data_[i]), x);
atomic_add<X>(c_style_pointer_cast<X*>(&p_data_[i]), x);
}
}
}