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

update build

Browse Source 
[ROCm/composable_kernel commit: a414e3fdf8]
This commit is contained in:
Chao Liu
2019-02-15 02:06:34 -06:00
parent e7f6b820cd
commit 153629655f
6 changed files with 17 additions and 11 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
driver/CMakeLists.txt
View File

@@ -1,4 +1,8 @@
set(DRIVER_SOURCE driver.cpp)
if(DEVICE_BACKEND STREQUAL "HIP")
set(DRIVER_SOURCE driver.hip.cpp)
elseif(DEVICE_BACKEND STREQUAL "CUDA")
set(DRIVER_SOURCE driver.cu)
endif()
add_executable(driver ${DRIVER_SOURCE})
target_link_libraries(driver PRIVATE tensor)

1
driver/driver.cu Symbolic link
View File

@@ -0,0 +1 @@
driver.hip.cpp

14
driver/driver.cpp → driver/driver.hip.cpp
View File

@@ -50,7 +50,7 @@ struct GeneratorTensor_3
std::initializer_list<std::size_t> ids = {static_cast<std::size_t>(is)...};
std::vector<std::size_t> lens(sizeof...(Is), 100);
std::vector<std::size_t> strides(sizeof...(Is), 1);
std::partial_sum(lens.rbegin(), lens.rbegin() + (sizeof...(Is)-1), strides.rbegin() + 1);
std::partial_sum(lens.rbegin(), lens.rbegin() + (sizeof...(Is) - 1), strides.rbegin() + 1);
return std::inner_product(ids.begin(), ids.end(), strides.begin(), std::size_t(0)) + 1;
#endif
}
@@ -340,7 +340,7 @@ void host_winograd_3x3_convolution(
std::size_t ho = OutTileSizeH * y + j;
for(int i = 0; i < OutTileSizeW; ++i)
{
std::size_t wo = OutTileSizeW * x + i;
std::size_t wo = OutTileSizeW * x + i;
out(n, k, ho, wo) = out_hold(n, k, y, x, j, i);
}
}
@@ -393,13 +393,13 @@ int main()
constexpr unsigned WPad = 0;
#elif 0
// 3x3, 34x34
constexpr unsigned N = 64;
constexpr unsigned C = 256;
constexpr unsigned N = 64;
constexpr unsigned C = 256;
constexpr unsigned HI = 34;
constexpr unsigned WI = 34;
constexpr unsigned K = 64;
constexpr unsigned S = 3;
constexpr unsigned R = 3;
constexpr unsigned K = 64;
constexpr unsigned S = 3;
constexpr unsigned R = 3;
constexpr unsigned HPad = 0;
constexpr unsigned WPad = 0;