mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-18 12:00:07 +00:00
8ffa0717f9
* Use thread cluster descriptor and explicit M_K 2d descriptor to simply Blockwise Reduction
* Change by replacing ReduceDims by NumReduceDims as Device Reduce interface template parameter
* Rename the folder name for the pool2d and reduce examples
* Update to reduction test scripts
* Add Readme for pool2d_fwd and reduce_blockwise examples
* Add support for int8_t reduction (ADD/AVG, MIN/MAX/AMAX)
* Tiny fix in reduce profiler and tiny update in reduce testing scripts
* Tiny fix in testing script profile_reduce_no_index.sh
* Tiny fix in testing script profile_reduce_no_index.sh
* Add support for bfp16 reduction (using bhalf_t = ushort)
* Tiny fix in amd_buffer_addressing.hpp
* Tiny change in script/profile_reduce_with_index.sh
* Use AccDataType for Beta value and use element_wise::PassThrough
* Use type_convert for type converting in host layer reduction
* Renaming and refining in Reduction profiler/device layer/examples
* Renaming and refining in Reduction profiler/device layer/examples
* Renaming all NumReduceDims to NumReduceDim
* Fix the leaked type_convert in ThreadwiseTensorSliceTransfer_v2
* Update to testing scripts to add bf16 support
* added more static_assert
* Remove buggy tunable configurations defined in device_reduce_instance_xxx.hpp
* Add static_assert to give compile-time warning for incorrect thread slice-size/vector-size configurations
* minor change
* Refine and fix (in GetWorkspaceSizeInBytes of MultiBlockPartialReduce) to make int8 completely pass
* Tiny renaming in gridwise_2d_reduction_multiblock_partial_reduce.hpp
* Tiny fix in script/profile_reduce_no_index.sh
* Refine in DeviceReduce layer with regard to using NumInvariantDim/NumReduceDim or InvariantDims/ReduceDims
* Generic renaming in host reduction and DeviceReduce layer
* Add support for 4-d all dimension reduction in the profiler and add_device_reduce_xxx instances
* Use multi-thread and simplification for host Reduction implementation
* Add ctest for reduction
* Update to clarify the using of data init method in produce_reduce/example_reduce/test_reduce/
* Update to the reduce CTest executables to enable default testing behavior when no command argument
* Renaming
Co-authored-by: Jianfeng yan <jfyan008@gmail.com>
[ROCm/composable_kernel commit: 9a8ee8a39a]
61 lines
2.2 KiB
Markdown
61 lines
2.2 KiB
Markdown
# Instructions for ```reduce_blockwise``` Example
|
|
|
|
## Docker script
|
|
```bash
|
|
docker run \
|
|
-it \
|
|
--rm \
|
|
--privileged \
|
|
--group-add sudo \
|
|
-w /root/workspace \
|
|
-v ${PATH_TO_LOCAL_WORKSPACE}:/root/workspace \
|
|
rocm/tensorflow:rocm4.3.1-tf2.6-dev \
|
|
/bin/bash
|
|
```
|
|
|
|
## Build ```reduce_blockwise```
|
|
```bash
|
|
mkdir build && cd build
|
|
```
|
|
|
|
```bash
|
|
# Need to specify target ID, example below is gfx908
|
|
cmake \
|
|
-D BUILD_DEV=OFF \
|
|
-D CMAKE_BUILD_TYPE=Release \
|
|
-D CMAKE_CXX_FLAGS="-DCK_AMD_GPU_GFX908 --amdgpu-target=gfx908 -O3 " \
|
|
-D CMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc \
|
|
-D CMAKE_PREFIX_PATH=/opt/rocm \
|
|
..
|
|
```
|
|
|
|
```bash
|
|
make -j reduce_blockwise
|
|
```
|
|
|
|
## Run ```reduce_blockwise```
|
|
```bash
|
|
# -D <xxx> : input 4-d tensor lengths
|
|
# -v <x> : verification (0=no, 1=yes)
|
|
#arg1: initialization (0=no init, 1=single integer value, 2=scope integer value, 3=decimal value)
|
|
#arg2: run kernel # of times (>1)
|
|
./bin/reduce_blockwise -D 16,64,32,960 -v 1 1 10
|
|
```
|
|
|
|
Result
|
|
```
|
|
launch_and_time_kernel: grid_dim {240, 1, 1}, block_dim {256, 1, 1}
|
|
Warm up
|
|
Start running 3 times...
|
|
Perf: 0.23536 ms, 267.32 GB/s, DeviceReduceBlockWise<256,M_C4_S1,K_C64_S1,InSrcVectorDim_0_InSrcVectorSize_1_OutDstVectorSize_1>
|
|
error: 0
|
|
max_diff: 0, 529, 529
|
|
root@dc-smc-18:/data/composable_kernel/Build3# bin/reduce_blockwise -D 16,64,32,960 -v 1 1 10
|
|
launch_and_time_kernel: grid_dim {240, 1, 1}, block_dim {256, 1, 1}
|
|
Warm up
|
|
Start running 10 times...
|
|
Perf: 0.23392 ms, 268.966 GB/s, DeviceReduceBlockWise<256,M_C4_S1,K_C64_S1,InSrcVectorDim_0_InSrcVectorSize_1_OutDstVectorSize_1>
|
|
error: 0
|
|
max_diff: 0, 528, 528
|
|
```
|