@@ -315,9 +315,9 @@ struct ExampleRunner {
315
315
SrcT* h_src = new SrcT[size * L];
316
316
ElementScale* scale_h = new ElementScale[L];
317
317
ElementZero* zero_h = new ElementZero[L];
318
- syclcompat ::memcpy (h_src, d_src, size * L * sizeof (SrcT));
319
- syclcompat ::memcpy (scale_h, scale, L * sizeof (ElementScale));
320
- syclcompat ::memcpy (zero_h, zero, L * sizeof (ElementZero));
318
+ cutlasscompat ::memcpy (h_src, d_src, size * L * sizeof (SrcT));
319
+ cutlasscompat ::memcpy (scale_h, scale, L * sizeof (ElementScale));
320
+ cutlasscompat ::memcpy (zero_h, zero, L * sizeof (ElementZero));
321
321
322
322
DstT* h_dst = new DstT[size * L];
323
323
for (size_t j = 0 ; j < L; ++j) {
@@ -326,7 +326,7 @@ struct ExampleRunner {
326
326
}
327
327
}
328
328
329
- syclcompat ::memcpy (d_dst, h_dst, size * sizeof (DstT));
329
+ cutlasscompat ::memcpy (d_dst, h_dst, size * sizeof (DstT));
330
330
}
331
331
332
332
// / Populates a Gemm::Arguments structure from the given commandline options
@@ -466,10 +466,10 @@ struct ExampleRunner {
466
466
CUTLASS_CHECK (gemm_ref.initialize (arguments, workspace.get ()));
467
467
CUTLASS_CHECK (gemm_ref.run ());
468
468
469
- syclcompat ::wait ();
469
+ cutlasscompat ::wait ();
470
470
// compare_reference
471
471
passed |= cutlass::reference::device::BlockCompareRelativelyEqual (block_ref_D.get (), block_D.get () + offset_D[i], block_ref_D.size (), epsilon, non_zero_floor);
472
- syclcompat ::wait ();
472
+ cutlasscompat ::wait ();
473
473
}
474
474
475
475
return passed;
@@ -615,7 +615,7 @@ struct ExampleRunner {
615
615
std::vector<uint8_t > zero (size (zero_layout) * sizeof_bits_v<ElementZero> / 8 , 0 );
616
616
cutlass::device_memory::copy_to_host (zero.data (), (uint8_t *)zero_buffer, zero.size ());
617
617
618
- syclcompat ::wait ();
618
+ cutlasscompat ::wait ();
619
619
620
620
auto dst_tensor = make_tensor (make_gmem_ptr (reinterpret_cast <DequantizedElement*>(dst.data ())), operand_layout);
621
621
@@ -669,7 +669,7 @@ struct ExampleRunner {
669
669
}
670
670
671
671
cutlass::device_memory::copy_to_device (dq_buffer, (DequantizedElement*)(raw_pointer_cast (dst_tensor.data ())), dst_tensor.size ());
672
- syclcompat ::wait ();
672
+ cutlasscompat ::wait ();
673
673
}
674
674
675
675
@@ -872,7 +872,7 @@ struct ExampleRunner {
872
872
// Run the GEMM
873
873
CUTLASS_CHECK (gemm_op.run ());
874
874
875
- syclcompat ::wait ();
875
+ cutlasscompat ::wait ();
876
876
877
877
// Verify that the result is correct
878
878
bool passed = verify (options);
@@ -886,7 +886,7 @@ struct ExampleRunner {
886
886
for (int i = 0 ; i < options.iterations ; ++i) {
887
887
gemm_op.run ();
888
888
}
889
- syclcompat ::wait ();
889
+ cutlasscompat ::wait ();
890
890
891
891
float cute_time = timer.seconds () / options.iterations ;
892
892
double cute_average_time = double (cute_time) / double (options.iterations );
0 commit comments