Skip to content

Commit d683438

Browse files
TejaX-AlaghariTejaX-Alaghari
committed
Changed type comment
1 parent 722ed60 commit d683438

File tree

5 files changed

+18
-24
lines changed

5 files changed

+18
-24
lines changed

‎clang/examples/DPCT/Runtime/nvcuda$$wmma$$fill_fragment.cu

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,10 +1,9 @@
11
// Option: --use-experimental-features=matrix
22
#include <mma.h>
33

4-
__global__ void test() {
4+
__global__ void test(float val) {
55
// Start
66
nvcuda::wmma::fragment<nvcuda::wmma::accumulator, 16, 16, 16, float> acc_frag;
7-
nvcuda::wmma::fill_fragment(acc_frag /* type fragment */,
8-
1.0f /* type value */);
7+
nvcuda::wmma::fill_fragment(acc_frag, val /*float*/);
98
// End
109
}

‎clang/examples/DPCT/Runtime/nvcuda$$wmma$$load_matrix_sync.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -6,7 +6,7 @@ __global__ void test(half *a, int row, int col, int lda) {
66
nvcuda::wmma::fragment<nvcuda::wmma::matrix_a, 16, 16, 16, half,
77
nvcuda::wmma::row_major>
88
a_frag;
9-
nvcuda::wmma::load_matrix_sync(a_frag /* type fragment */,
10-
a + col + row * lda, lda);
9+
nvcuda::wmma::load_matrix_sync(a_frag, a + col + row * lda /*void **/,
10+
lda /*int*/);
1111
// End
1212
}

‎clang/examples/DPCT/Runtime/nvcuda$$wmma$$mma_sync.cu

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -10,8 +10,6 @@ __global__ void test() {
1010
nvcuda::wmma::col_major>
1111
b_frag;
1212
nvcuda::wmma::fragment<nvcuda::wmma::accumulator, 16, 16, 16, float> acc_frag;
13-
nvcuda::wmma::mma_sync(acc_frag /* type fragment */,
14-
a_frag /* type fragment */, b_frag /* type fragment */,
15-
acc_frag /* type fragment */);
13+
nvcuda::wmma::mma_sync(acc_frag, a_frag, b_frag, acc_frag);
1614
// End
1715
}

‎clang/examples/DPCT/Runtime/nvcuda$$wmma$$store_matrix_sync.cu

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -5,10 +5,10 @@ __global__ void test(float *c, int row, int col, int ldc) {
55
// Start
66
nvcuda::wmma::fragment<nvcuda::wmma::accumulator, 16, 16, 16, float> acc_frag;
77
nvcuda::wmma::store_matrix_sync(
8-
c + col + row * ldc, acc_frag /* type fragment */, ldc,
9-
nvcuda::wmma::mem_col_major /* type memory order */);
8+
c + col + row * ldc /*void **/, acc_frag, ldc /*int*/,
9+
nvcuda::wmma::mem_col_major /*memory order*/);
1010
nvcuda::wmma::store_matrix_sync(
11-
c + row + col * ldc, acc_frag /* type fragment */, ldc,
12-
nvcuda::wmma::mem_row_major /* type memory order */);
11+
c + row + col * ldc /*void **/, acc_frag, ldc /*int*/,
12+
nvcuda::wmma::mem_row_major /*memory order*/);
1313
// End
1414
}

‎clang/test/dpct/query_api_mapping/Runtime/test_wmma.cu

Lines changed: 9 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -4,19 +4,18 @@
44
// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=nvcuda::wmma::fill_fragment | FileCheck %s -check-prefix=NVCUDA_WMMA_FILL_FRAGMENT
55
// NVCUDA_WMMA_FILL_FRAGMENT: CUDA API:
66
// NVCUDA_WMMA_FILL_FRAGMENT-NEXT: nvcuda::wmma::fragment<nvcuda::wmma::accumulator, 16, 16, 16, float> acc_frag;
7-
// NVCUDA_WMMA_FILL_FRAGMENT-NEXT: nvcuda::wmma::fill_fragment(acc_frag /* type fragment */,
8-
// NVCUDA_WMMA_FILL_FRAGMENT-NEXT: 1.0f /* type value */);
7+
// NVCUDA_WMMA_FILL_FRAGMENT-NEXT: nvcuda::wmma::fill_fragment(acc_frag, val /*float*/);
98
// NVCUDA_WMMA_FILL_FRAGMENT-NEXT: Is migrated to (with the option --use-experimental-features=matrix):
109
// NVCUDA_WMMA_FILL_FRAGMENT-NEXT: dpct::experimental::matrix::joint_matrix<dpct::experimental::matrix::accumulator, 16, 16, 16, float> acc_frag;
11-
// NVCUDA_WMMA_FILL_FRAGMENT-NEXT: sycl::ext::oneapi::experimental::matrix::joint_matrix_fill(sycl::ext::oneapi::this_work_item::get_sub_group(), acc_frag.get(), 1.0f);
10+
// NVCUDA_WMMA_FILL_FRAGMENT-NEXT: sycl::ext::oneapi::experimental::matrix::joint_matrix_fill(sycl::ext::oneapi::this_work_item::get_sub_group(), acc_frag.get(), val);
1211

1312
// RUN: dpct --cuda-include-path="%cuda-path/include" --query-api-mapping=nvcuda::wmma::load_matrix_sync | FileCheck %s -check-prefix=NVCUDA_WMMA_LOAD_MATRIX_SYNC
1413
// NVCUDA_WMMA_LOAD_MATRIX_SYNC: CUDA API:
1514
// NVCUDA_WMMA_LOAD_MATRIX_SYNC-NEXT: nvcuda::wmma::fragment<nvcuda::wmma::matrix_a, 16, 16, 16, half,
1615
// NVCUDA_WMMA_LOAD_MATRIX_SYNC-NEXT: nvcuda::wmma::row_major>
1716
// NVCUDA_WMMA_LOAD_MATRIX_SYNC-NEXT: a_frag;
18-
// NVCUDA_WMMA_LOAD_MATRIX_SYNC-NEXT: nvcuda::wmma::load_matrix_sync(a_frag /* type fragment */,
19-
// NVCUDA_WMMA_LOAD_MATRIX_SYNC-NEXT: a + col + row * lda, lda);
17+
// NVCUDA_WMMA_LOAD_MATRIX_SYNC-NEXT: nvcuda::wmma::load_matrix_sync(a_frag, a + col + row * lda /*void **/,
18+
// NVCUDA_WMMA_LOAD_MATRIX_SYNC-NEXT: lda /*int*/);
2019
// NVCUDA_WMMA_LOAD_MATRIX_SYNC-NEXT: Is migrated to (with the option --use-experimental-features=matrix):
2120
// NVCUDA_WMMA_LOAD_MATRIX_SYNC-NEXT: dpct::experimental::matrix::joint_matrix<dpct::experimental::matrix::a, 16, 16, 16, sycl::half, dpct::experimental::matrix::row_major>
2221
// NVCUDA_WMMA_LOAD_MATRIX_SYNC-NEXT: a_frag;
@@ -26,11 +25,11 @@
2625
// NVCUDA_WMMA_STORE_MATRIX_SYNC: CUDA API:
2726
// NVCUDA_WMMA_STORE_MATRIX_SYNC-NEXT: nvcuda::wmma::fragment<nvcuda::wmma::accumulator, 16, 16, 16, float> acc_frag;
2827
// NVCUDA_WMMA_STORE_MATRIX_SYNC-NEXT: nvcuda::wmma::store_matrix_sync(
29-
// NVCUDA_WMMA_STORE_MATRIX_SYNC-NEXT: c + col + row * ldc, acc_frag /* type fragment */, ldc,
30-
// NVCUDA_WMMA_STORE_MATRIX_SYNC-NEXT: nvcuda::wmma::mem_col_major /* type memory order */);
28+
// NVCUDA_WMMA_STORE_MATRIX_SYNC-NEXT: c + col + row * ldc /*void **/, acc_frag, ldc /*int*/,
29+
// NVCUDA_WMMA_STORE_MATRIX_SYNC-NEXT: nvcuda::wmma::mem_col_major /*memory order*/);
3130
// NVCUDA_WMMA_STORE_MATRIX_SYNC-NEXT: nvcuda::wmma::store_matrix_sync(
32-
// NVCUDA_WMMA_STORE_MATRIX_SYNC-NEXT: c + row + col * ldc, acc_frag /* type fragment */, ldc,
33-
// NVCUDA_WMMA_STORE_MATRIX_SYNC-NEXT: nvcuda::wmma::mem_row_major /* type memory order */);
31+
// NVCUDA_WMMA_STORE_MATRIX_SYNC-NEXT: c + row + col * ldc /*void **/, acc_frag, ldc /*int*/,
32+
// NVCUDA_WMMA_STORE_MATRIX_SYNC-NEXT: nvcuda::wmma::mem_row_major /*memory order*/);
3433
// NVCUDA_WMMA_STORE_MATRIX_SYNC-NEXT: Is migrated to (with the option --use-experimental-features=matrix):
3534
// NVCUDA_WMMA_STORE_MATRIX_SYNC-NEXT: dpct::experimental::matrix::joint_matrix<dpct::experimental::matrix::accumulator, 16, 16, 16, float> acc_frag;
3635
// NVCUDA_WMMA_STORE_MATRIX_SYNC-NEXT: sycl::ext::oneapi::experimental::matrix::joint_matrix_store(sycl::ext::oneapi::this_work_item::get_sub_group(), acc_frag.get(), sycl::address_space_cast<sycl::access::address_space::generic_space, sycl::access::decorated::no, float>(c + col + row * ldc), ldc, sycl::ext::oneapi::experimental::matrix::layout::col_major);
@@ -45,9 +44,7 @@
4544
// NVCUDA_WMMA_MMA_SYNC-NEXT: nvcuda::wmma::col_major>
4645
// NVCUDA_WMMA_MMA_SYNC-NEXT: b_frag;
4746
// NVCUDA_WMMA_MMA_SYNC-NEXT: nvcuda::wmma::fragment<nvcuda::wmma::accumulator, 16, 16, 16, float> acc_frag;
48-
// NVCUDA_WMMA_MMA_SYNC-NEXT: nvcuda::wmma::mma_sync(acc_frag /* type fragment */,
49-
// NVCUDA_WMMA_MMA_SYNC-NEXT: a_frag /* type fragment */, b_frag /* type fragment */,
50-
// NVCUDA_WMMA_MMA_SYNC-NEXT: acc_frag /* type fragment */);
47+
// NVCUDA_WMMA_MMA_SYNC-NEXT: nvcuda::wmma::mma_sync(acc_frag, a_frag, b_frag, acc_frag);
5148
// NVCUDA_WMMA_MMA_SYNC-NEXT: Is migrated to (with the option --use-experimental-features=matrix):
5249
// NVCUDA_WMMA_MMA_SYNC-NEXT: dpct::experimental::matrix::joint_matrix<dpct::experimental::matrix::a, 16, 16, 16, sycl::half, dpct::experimental::matrix::row_major>
5350
// NVCUDA_WMMA_MMA_SYNC-NEXT: a_frag;

0 commit comments

Comments
 (0)