Skip to content
GitLab
Menu
Projects
Groups
Snippets
Loading...
Help
Help
Support
Community forum
Keyboard shortcuts
?
Submit feedback
Contribute to GitLab
Sign in / Register
Toggle navigation
Menu
Open sidebar
gaoqiong
composable_kernel
Commits
b42fe7c3
Commit
b42fe7c3
authored
Oct 13, 2023
by
Bartlomiej Kocot
Browse files
Minor fixes
parent
92d725cf
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
3 additions
and
4 deletions
+3
-4
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_utils.hpp
...r_operation/gpu/device/impl/device_grouped_conv_utils.hpp
+2
-3
test/grouped_convnd_bwd_weight/CMakeLists.txt
test/grouped_convnd_bwd_weight/CMakeLists.txt
+1
-1
No files found.
include/ck/tensor_operation/gpu/device/impl/device_grouped_conv_utils.hpp
View file @
b42fe7c3
...
@@ -43,14 +43,13 @@ struct ComputePtrOffsetOfStridedBatch
...
@@ -43,14 +43,13 @@ struct ComputePtrOffsetOfStridedBatch
return
ds_offset
;
return
ds_offset
;
}
}
__host__
__device__
constexpr
long_index_t
GetEPtrOffset
(
index_t
g_idx
)
const
[[
maybe_unused
]]
__host__
__device__
constexpr
long_index_t
GetEPtrOffset
(
index_t
g_idx
)
const
{
{
return
g_idx
*
static_cast
<
long_index_t
>
(
BatchStrideE_
);
return
g_idx
*
static_cast
<
long_index_t
>
(
BatchStrideE_
);
}
}
// alias for kernels without multiple D
// alias for kernels without multiple D
template
<
typename
enable_if
<
NumDTensor
==
0
,
bool
>
::
type
=
false
>
[[
maybe_unused
]]
__host__
__device__
constexpr
long_index_t
GetCPtrOffset
(
index_t
g_idx
)
const
__host__
__device__
constexpr
long_index_t
GetCPtrOffset
(
index_t
g_idx
)
const
{
{
return
g_idx
*
static_cast
<
long_index_t
>
(
BatchStrideE_
);
return
g_idx
*
static_cast
<
long_index_t
>
(
BatchStrideE_
);
}
}
...
...
test/grouped_convnd_bwd_weight/CMakeLists.txt
View file @
b42fe7c3
...
@@ -12,7 +12,7 @@ foreach(gpu IN LISTS GPU_TARGETS)
...
@@ -12,7 +12,7 @@ foreach(gpu IN LISTS GPU_TARGETS)
endif
()
endif
()
if
(
gpu IN_LIST gpu_list_wmma AND target EQUAL 0
)
if
(
gpu IN_LIST gpu_list_wmma AND target EQUAL 0
)
add_gtest_executable
(
test_grouped_convnd_bwd_weight test_grouped_convnd_bwd_weight.cpp
)
add_gtest_executable
(
test_grouped_convnd_bwd_weight test_grouped_convnd_bwd_weight.cpp
)
target_link_libraries
(
test_grouped_convnd_bwd_weight PRIVATE utility device_grouped_conv3d_bwd_weight_instance
)
target_link_libraries
(
test_grouped_convnd_bwd_weight PRIVATE utility
device_grouped_conv1d_bwd_weight_instance device_grouped_conv2d_bwd_weight_instance
device_grouped_conv3d_bwd_weight_instance
)
add_gtest_executable
(
test_grouped_convnd_bwd_weight_interface test_grouped_convnd_bwd_weight_interface_wmma.cpp
)
add_gtest_executable
(
test_grouped_convnd_bwd_weight_interface test_grouped_convnd_bwd_weight_interface_wmma.cpp
)
target_link_libraries
(
test_grouped_convnd_bwd_weight_interface PRIVATE utility
)
target_link_libraries
(
test_grouped_convnd_bwd_weight_interface PRIVATE utility
)
set
(
target 1
)
set
(
target 1
)
...
...
Write
Preview
Markdown
is supported
0%
Try again
or
attach a new file
.
Attach a file
Cancel
You are about to add
0
people
to the discussion. Proceed with caution.
Finish editing this message first!
Cancel
Please
register
or
sign in
to comment