onednn-src 0.1.13

Source of oneAPI Deep Neural Network Library (oneDNN)
Documentation 
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188

/*******************************************************************************
* Copyright 2020 Intel Corporation
* Copyright 2020 Codeplay Software Limited
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
*     http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*******************************************************************************/

#ifndef GPU_AMD_MIOPEN_INNER_PRODUCT_IMPL_HPP
#define GPU_AMD_MIOPEN_INNER_PRODUCT_IMPL_HPP

#include "rocblas/rocblas.h"
#include <miopen/miopen.h>

#include "common/type_helpers.hpp"
#include "gpu/amd/engine.hpp"
#include "gpu/amd/sycl_hip_utils.hpp"

namespace dnnl {
namespace impl {
namespace gpu {
namespace amd {
namespace {
inline void get_4d_tensor_descriptor(
        const memory_desc_t *mem_desc1, int *dims, int *strides) {
    memory_desc_t mem_desc = *mem_desc1;

    // Forcing tensors dims less than 4 to be 4 {n c h w};
    using namespace format_tag;
    auto set_dim = [&]() {
        if (mem_desc.ndims == 3) {
            mem_desc.ndims = 4;
            mem_desc.dims[3] = mem_desc.dims[2];
            mem_desc.dims[2] = 1;
            mem_desc.padded_dims[3] = mem_desc.padded_dims[2];
            mem_desc.padded_dims[2] = 1;
        } else if (mem_desc.ndims == 2) {
            mem_desc.ndims = 4;
            mem_desc.dims[3] = 1;
            mem_desc.dims[2] = 1;
            mem_desc.padded_dims[3] = 1;
            mem_desc.padded_dims[2] = 1;
        }
    };
    // Forcing strides < 4 to be 4
    if (memory_desc_matches_tag(mem_desc, nwc)) {
        set_dim();
        //  promoting nwc(owi) to NHWC = {wc 1 c} to {wc 1 wc c}
        mem_desc.format_desc.blocking.strides[3]
                = mem_desc.format_desc.blocking.strides[2];
        mem_desc.format_desc.blocking.strides[2]
                = mem_desc.format_desc.blocking.strides[0];
        assert(memory_desc_matches_tag(mem_desc, nhwc)
                && "Tag is not set to NHWC");
    } else if (memory_desc_matches_tag(mem_desc, ncw)) {
        set_dim();
        // promoting ncw(oiw) to NCHW = {wc w 1} to {wc w w 1}
        mem_desc.format_desc.blocking.strides[3]
                = mem_desc.format_desc.blocking.strides[2];
        mem_desc.format_desc.blocking.strides[2]
                = mem_desc.format_desc.blocking.strides[1];
        assert(memory_desc_matches_tag(mem_desc, nchw)
                && "Tag is not set to NCHW");
    } else if (memory_desc_matches_tag(mem_desc, wio)) {
        set_dim();
        // promoting wcn(wio) to HWCN = {1 n nc} to {1 n ncw nc}
        mem_desc.format_desc.blocking.strides[3]
                = mem_desc.format_desc.blocking.strides[2];
        mem_desc.format_desc.blocking.strides[2] *= mem_desc.dims[3];
        assert(memory_desc_matches_tag(mem_desc, hwio)
                && " Tag is not set to HWIO");
    } else if (memory_desc_matches_tag(mem_desc, nc)) {
        set_dim();
        // fixing strides
        // promoting nc(oi) to NCHW = {c 1} to {c 1 1 1}
        mem_desc.format_desc.blocking.strides[2]
                = mem_desc.format_desc.blocking.strides[1];
        mem_desc.format_desc.blocking.strides[3]
                = mem_desc.format_desc.blocking.strides[1];
        assert(memory_desc_matches_tag(mem_desc, nchw)
                && " Tag is not set to NCHW");
    } else if (memory_desc_matches_tag(mem_desc, cn)) {
        set_dim();
        // fixing strides cn(oi) to HWCN = {1 n} to {1 n nc nc}.
        // Note that CHWN exists as well, but for inner product
        // we convert it to HWCN. Other primitives may need
        // different conversion.
        mem_desc.format_desc.blocking.strides[2]
                = mem_desc.format_desc.blocking.strides[1]
                * mem_desc.padded_dims[1];
        mem_desc.format_desc.blocking.strides[3]
                = mem_desc.format_desc.blocking.strides[2];
        assert(memory_desc_matches_tag(mem_desc, hwio)
                && " Tag is not set to NCHW");
    }
    convert_dnnl_dims_array(mem_desc.dims, dims, mem_desc.ndims);
    convert_dnnl_dims_array(
            mem_desc.format_desc.blocking.strides, strides, mem_desc.ndims);
}
} // namespace
struct miopen_inner_product_impl_base_t {
    // The io enum requires the weights be the last parameter to ensure
    // tensor_descs is contiguous.
    enum io { src = 0, bia, dst, wei, NUM_IO };
    miopenDataType_t data_types_[NUM_IO + 1]; // +1 data-type for accumulation
    int ndims_;
    int dims_[NUM_IO][DNNL_MAX_NDIMS];
    // one extra stride added for transform filter
    int strides_[NUM_IO + 1][DNNL_MAX_NDIMS];

    miopenTensorDescriptor_t tensor_descs_[NUM_IO - 1] = {};

    size_t workspace_size_ = 0;
    float alpha_ = 1, beta_ = 0;
    bool with_bias_;
    bool scale_bias_ = false;
    bool with_relu_ = false, with_eltwise_ = false, with_sum_ = false;
    bool filter_using_spatial_format_ = false;

    virtual bool need_to_transform_filter() const {
        return filter_using_spatial_format_;
    }

    virtual bool ip_using_scratchpad() const { return (workspace_size_ > 0); }
    bool conv_using_scale_scratchpad() const { return scale_bias_; }

    void set_bias_dims(miopenTensorLayout_t format, int ndims, int bias_dim) {
        for (int i = 0; i < ndims; ++i) {
            dims_[io::bia][i] = 1;
            strides_[io::bia][i] = (format != miopenTensorNHWC ? 1 : bias_dim);
        }
        dims_[io::bia][1] = bias_dim;
        strides_[io::bia][1] = 1;
        strides_[io::bia][0] = bias_dim;
    }
    virtual status_t init(impl::engine_t * /*engine*/,
            inner_product_pd_t * /*pd*/, bool /*with_relu*/,
            bool /*with_eltwise*/, bool /*with_sum */,
            bool /*using_fused_path_for_blocking*/)
            = 0;

    virtual void execute(miopenHandle_t /*handle*/,
            rocblas_handle /*rocblas_handle*/,
            const std::vector<void *> & /*args*/) const
            = 0;
};

struct miopen_inner_product_fwd_base_t
    : public miopen_inner_product_impl_base_t {
    float alpha_; // alpha in gemm
    bool do_scaling_ {false}, runtime_scaling_ {false};
    float sum_scale_; // beta in gemm
    float eltwise_alpha(const inner_product_pd_t *pd) const {
        const int eltwise_idx
                = pd->attr()->post_ops_.find(primitive_kind::eltwise);
        return with_eltwise_
                ? pd->attr()->post_ops_.entry_[eltwise_idx].eltwise.alpha
                : 0.0f;
    }
    float sum_scale(const inner_product_pd_t *pd) const {
        const int sum_idx = pd->attr()->post_ops_.find(primitive_kind::sum);
        return with_sum_ ? pd->attr()->post_ops_.entry_[sum_idx].sum.scale
                         : 0.0f;
    }

    dnnl::impl::alg_kind_t eltwise_algorithm_kind(
            const inner_product_pd_t *pd) const {
        const int eltwise_idx
                = pd->attr()->post_ops_.find(primitive_kind::eltwise);
        return pd->attr()->post_ops_.entry_[eltwise_idx].eltwise.alg;
    }
};

} // namespace amd
} // namespace gpu
} // namespace impl
} // namespace dnnl

#endif