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
MIGraphX
Commits
2593dd60
"git@developer.sourcefind.cn:orangecat/ollama.git" did not exist on "7194a07d4d2e896e397c1e49d93b6a5fc2009972"
Commit
2593dd60
authored
Sep 13, 2022
by
turneram
Browse files
Formatting
parent
d1e27426
Changes
2
Hide whitespace changes
Inline
Side-by-side
Showing
2 changed files
with
8 additions
and
8 deletions
+8
-8
src/targets/gpu/jit/ck_gemm.cpp
src/targets/gpu/jit/ck_gemm.cpp
+7
-7
src/targets/gpu/kernels/include/migraphx/kernels/ck_gemm.hpp
src/targets/gpu/kernels/include/migraphx/kernels/ck_gemm.hpp
+1
-1
No files found.
src/targets/gpu/jit/ck_gemm.cpp
View file @
2593dd60
...
@@ -53,18 +53,18 @@ namespace gpu {
...
@@ -53,18 +53,18 @@ namespace gpu {
// extern "C" {
// extern "C" {
// __global__ void ck_gemm_kernel(void* a_p, void* b_p, void* c_p)
// __global__ void ck_gemm_kernel(void* a_p, void* b_p, void* c_p)
// {
// {
// // hipDeviceProp_t hdp{};
// // hipDeviceProp_t hdp{};
// // printf("Shared mem: %i\n", int(hdp.sharedMemPerBlock));
// // printf("Shared mem: %i\n", int(hdp.sharedMemPerBlock));
// // make_tensors()(a_p, b_p, c_p)([](auto&&... xs) {
// // make_tensors()(a_p, b_p, c_p)([](auto&&... xs) {
// // ck_gemm(xs...);
// // ck_gemm(xs...);
// // });
// // });
// make_tensors()(a_p, b_p, c_p)([](auto a_t, auto b_t, auto c_t) {
// make_tensors()(a_p, b_p, c_p)([](auto a_t, auto b_t, auto c_t) {
// __shared__ float p_shared_block[512]; //[(a_t.get_shape().elements() +
b_t.get_shape().elements()) * 2];
// __shared__ float p_shared_block[512]; //[(a_t.get_shape().elements() +
// ck_gemm(a_t, b_t, c_t, p_shared_block);
//
b_t.get_shape().elements()) * 2];
ck_gemm(a_t, b_t, c_t, p_shared_block);
// // make_tensors()(p_shared_block)([&](auto p_t) {
// // make_tensors()(p_shared_block)([&](auto p_t) {
// // ck_gemm(a_t, b_t, c_t, p_t);
// // ck_gemm(a_t, b_t, c_t, p_t);
// // });
// // });
// });
// });
// }
// }
...
...
src/targets/gpu/kernels/include/migraphx/kernels/ck_gemm.hpp
View file @
2593dd60
...
@@ -283,7 +283,7 @@ __device__ void ck_gemm(const T& a_t, const U& b_t, const V& c_t, float* p_t)
...
@@ -283,7 +283,7 @@ __device__ void ck_gemm(const T& a_t, const U& b_t, const V& c_t, float* p_t)
constexpr
bool
HasMainKBlockLoop
=
true
;
constexpr
bool
HasMainKBlockLoop
=
true
;
constexpr
bool
HasDoubleTailKBlockLoop
=
true
;
constexpr
bool
HasDoubleTailKBlockLoop
=
true
;
auto
num_bytes
=
GridwiseGemm
::
GetSharedMemoryNumberOfByte
();
auto
num_bytes
=
GridwiseGemm
::
GetSharedMemoryNumberOfByte
();
printf
(
"Bytes: %i
\n
"
,
int
(
num_bytes
));
printf
(
"Bytes: %i
\n
"
,
int
(
num_bytes
));
GridwiseGemm
::
Run
(
a_t
.
data
(),
GridwiseGemm
::
Run
(
a_t
.
data
(),
b_t
.
data
(),
b_t
.
data
(),
...
...
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