Repository navigation
Expand file tree
/
Copy pathwgrads.cpp
More file actions
134 lines (104 loc) · 5.78 KB
/
Copy pathwgrads.cpp
File metadata and controls
134 lines (104 loc) · 5.78 KB
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
/*
* SPDX-FileCopyrightText: Copyright (c) 2023 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
* SPDX-License-Identifier: MIT
*/
#include <catch2/catch_test_macros.hpp>
#include "../utils/helpers.h"
#include <cudnn_frontend.h>
TEST_CASE("Convolution Wgrad", "[wgrad][graph][wgrad][Conv_wgrad]") {
namespace fe = cudnn_frontend;
if (is_arch_supported_by_cudnn() == false) {
SKIP("Architecture is not supported by currend cudnn version");
}
fe::graph::Graph graph;
graph.set_io_data_type(fe::DataType_t::HALF)
.set_intermediate_data_type(fe::DataType_t::HALF)
.set_compute_data_type(fe::DataType_t::FLOAT);
auto X = graph.tensor(fe::graph::Tensor_attributes()
.set_name("image")
.set_dim({4, 64, 16, 16})
.set_stride({64 * 16 * 16, 1, 64 * 16, 64}));
auto DY = graph.tensor(fe::graph::Tensor_attributes()
.set_name("grad")
.set_dim({4, 64, 16, 16})
.set_stride({64 * 16 * 16, 1, 64 * 16, 64}));
auto wgrad_options = fe::graph::Conv_wgrad_attributes().set_padding({1, 1}).set_stride({1, 1}).set_dilation({1, 1});
auto DW = graph.conv_wgrad(DY, X, wgrad_options);
DW->set_output(true).set_dim({64, 64, 3, 3});
// Create a unique_ptr for the cuDNN handle
auto handle_ptr = create_cudnn_handle();
auto handle = *handle_ptr;
REQUIRE(graph.validate().is_good());
REQUIRE(graph.build_operation_graph(handle).is_good());
REQUIRE(graph.create_execution_plans({fe::HeurMode_t::A}).is_good());
REQUIRE(graph.check_support().is_good());
REQUIRE(graph.build_plans().is_good());
Surface<half> x_tensor(4 * 64 * 16 * 16);
Surface<half> dy_tensor(4 * 64 * 16 * 16);
Surface<half> dw_tensor(64 * 64 * 3 * 3);
int64_t workspace_size = 0;
REQUIRE(graph.get_workspace_size(workspace_size).is_good());
Surface<int8_t> workspace(workspace_size);
std::unordered_map<std::shared_ptr<fe::graph::Tensor_attributes>, void*> variant_pack = {
{X, x_tensor.devPtr}, {DY, dy_tensor.devPtr}, {DW, dw_tensor.devPtr}};
REQUIRE(graph.execute(handle, variant_pack, workspace.devPtr).is_good());
}
TEST_CASE("scale-bias-relu-wgrad Graph", "[wgrad][graph][scale-bias-relu-wgrad][ConvBNwgrad]") {
if (!is_ampere_arch() && !is_hopper_arch()) {
SKIP("scale-bias-relu-wgrad requires Ampere or Hopper");
}
namespace fe = cudnn_frontend;
fe::graph::Graph graph;
graph.set_io_data_type(fe::DataType_t::HALF)
.set_intermediate_data_type(fe::DataType_t::HALF)
.set_compute_data_type(fe::DataType_t::FLOAT);
auto X = graph.tensor(fe::graph::Tensor_attributes()
.set_name("image")
.set_dim({4, 64, 16, 16})
.set_stride({64 * 16 * 16, 1, 64 * 16, 64}));
auto S = graph.tensor(
fe::graph::Tensor_attributes().set_name("scale").set_dim({1, 64, 1, 1}).set_stride({64, 1, 64, 64}));
auto scale_options = fe::graph::Pointwise_attributes().set_mode(fe::PointwiseMode_t::MUL);
auto scale_output = graph.pointwise(X, S, scale_options);
auto B = graph.tensor(
fe::graph::Tensor_attributes().set_name("bias").set_dim({1, 64, 1, 1}).set_stride({64, 1, 64, 64}));
auto bias_options = fe::graph::Pointwise_attributes().set_mode(fe::PointwiseMode_t::ADD);
auto bias_output = graph.pointwise(scale_output, B, bias_options);
auto relu_options = fe::graph::Pointwise_attributes().set_mode(fe::PointwiseMode_t::RELU_FWD);
auto relu_output = graph.pointwise(bias_output, relu_options);
auto DY = graph.tensor(fe::graph::Tensor_attributes()
.set_name("grad")
.set_dim({4, 64, 16, 16})
.set_stride({64 * 16 * 16, 1, 64 * 16, 64}));
auto wgrad_options = fe::graph::Conv_wgrad_attributes().set_padding({1, 1}).set_stride({1, 1}).set_dilation({1, 1});
auto DW = graph.conv_wgrad(DY, relu_output, wgrad_options);
DW->set_output(true).set_dim({64, 64, 3, 3});
#if (CUDNN_VERSION < 8800)
SKIP("ConvBNwgrad requires cudnn 8.8 and up");
#endif
if (check_device_arch_newer_than("ampere") == false) {
SKIP("ConvBNwgrad requires hopper and above architecture.");
}
// Create a unique_ptr for the cuDNN handle
auto handle_ptr = create_cudnn_handle();
auto handle = *handle_ptr;
REQUIRE(graph.validate().is_good());
REQUIRE(graph.build_operation_graph(handle).is_good());
REQUIRE(graph.create_execution_plans({fe::HeurMode_t::A}).is_good());
REQUIRE(graph.check_support().is_good());
REQUIRE(graph.build_plans().is_good());
Surface<half> x_tensor(4 * 64 * 16 * 16);
Surface<half> s_tensor(64);
Surface<half> b_tensor(64);
Surface<half> dy_tensor(4 * 64 * 16 * 16);
Surface<half> dw_tensor(64 * 64 * 3 * 3);
int64_t workspace_size = 0;
REQUIRE(graph.get_workspace_size(workspace_size).is_good());
Surface<int8_t> workspace(workspace_size);
std::unordered_map<std::shared_ptr<fe::graph::Tensor_attributes>, void*> variant_pack = {{X, x_tensor.devPtr},
{S, s_tensor.devPtr},
{B, b_tensor.devPtr},
{DY, dy_tensor.devPtr},
{DW, dw_tensor.devPtr}};
REQUIRE(graph.execute(handle, variant_pack, workspace.devPtr).is_good());
}