AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Merge branch 'master' of https://github.com/ROCm-Developer-Tools/HIP into feature_purge_genco

Browse Source 
[ROCm/hip commit: ed75522ba5]
This commit is contained in:
Alex Voicu
2018-01-17 14:02:19 +00:00
parent 044ae14b92 5e86facde3
commit abb22cc7f2
35 changed files with 529 additions and 413 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/tests/hipify-clang/axpy.cu
-1
View File
@@ -16,7 +16,6 @@
template<typename T>
__global__ void axpy(T a, T *x, T *y) {
// CHECK: y[hipThreadIdx_x] = a * x[hipThreadIdx_x];
y[threadIdx.x] = a * x[threadIdx.x];
}
projects/hip/tests/hipify-clang/concurentKernels.cu
-242
View File
@@ -1,242 +0,0 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
/*
* Copyright 1993-2015 NVIDIA Corporation. All rights reserved.
*
* Please refer to the NVIDIA end user license agreement (EULA) associated
* with this source code for terms and conditions that govern your use of
* this software. Any use, reproduction, disclosure, or distribution of
* this software and related documentation outside the terms of the EULA
* is strictly prohibited.
*
*/
//
// This sample demonstrates the use of streams for concurrent execution. It also illustrates how to
// introduce dependencies between CUDA streams with the new cudaStreamWaitEvent function introduced
// in CUDA 3.2.
//
// Devices of compute capability 1.x will run the kernels one after another
// Devices of compute capability 2.0 or higher can overlap the kernels
//
#include <stdio.h>
#include <helper_functions.h>
#include <helper_cuda.h>
// This is a kernel that does no real work but runs at least for a specified number of clocks
__global__ void clock_block(clock_t *d_o, clock_t clock_count)
{
unsigned int start_clock = (unsigned int) clock();
clock_t clock_offset = 0;
while (clock_offset < clock_count)
{
unsigned int end_clock = (unsigned int) clock();
// The code below should work like
// this (thanks to modular arithmetics):
//
// clock_offset = (clock_t) (end_clock > start_clock ?
// end_clock - start_clock :
// end_clock + (0xffffffffu - start_clock));
//
// Indeed, let m = 2^32 then
// end - start = end + m - start (mod m).
clock_offset = (clock_t)(end_clock - start_clock);
}
d_o[0] = clock_offset;
}
// Single warp reduction kernel
__global__ void sum(clock_t *d_clocks, int N)
{
__shared__ clock_t s_clocks[32];
clock_t my_sum = 0;
for (int i = threadIdx.x; i < N; i+= blockDim.x)
{
my_sum += d_clocks[i];
}
s_clocks[threadIdx.x] = my_sum;
syncthreads();
for (int i=16; i>0; i/=2)
{
if (threadIdx.x < i)
{
s_clocks[threadIdx.x] += s_clocks[threadIdx.x + i];
}
syncthreads();
}
d_clocks[0] = s_clocks[0];
}
int main(int argc, char **argv)
{
int nkernels = 8; // number of concurrent kernels
int nstreams = nkernels + 1; // use one more stream than concurrent kernel
int nbytes = nkernels * sizeof(clock_t); // number of data bytes
float kernel_time = 10; // time the kernel should run in ms
float elapsed_time; // timing variables
int cuda_device = 0;
printf("[%s] - Starting...\n", argv[0]);
// get number of kernels if overridden on the command line
if (checkCmdLineFlag(argc, (const char **)argv, "nkernels"))
{
nkernels = getCmdLineArgumentInt(argc, (const char **)argv, "nkernels");
nstreams = nkernels + 1;
}
// use command-line specified CUDA device, otherwise use device with highest Gflops/s
cuda_device = findCudaDevice(argc, (const char **)argv);
// CHECK: hipDeviceProp_t deviceProp;
cudaDeviceProp deviceProp;
// CHECK: checkCudaErrors(hipGetDevice(&cuda_device));
checkCudaErrors(cudaGetDevice(&cuda_device));
// CHECK: checkCudaErrors(hipGetDeviceProperties(&deviceProp, cuda_device));
checkCudaErrors(cudaGetDeviceProperties(&deviceProp, cuda_device));
if ((deviceProp.concurrentKernels == 0))
{
printf("> GPU does not support concurrent kernel execution\n");
printf(" CUDA kernel runs will be serialized\n");
}
printf("> Detected Compute SM %d.%d hardware with %d multi-processors\n",
deviceProp.major, deviceProp.minor, deviceProp.multiProcessorCount);
// allocate host memory
clock_t *a = 0; // pointer to the array data in host memory
// CHECK: checkCudaErrors(hipHostMalloc((void **)&a, nbytes));
checkCudaErrors(cudaMallocHost((void **)&a, nbytes));
// allocate device memory
clock_t *d_a = 0; // pointers to data and init value in the device memory
// CHECK: checkCudaErrors(hipMalloc((void **)&d_a, nbytes));
checkCudaErrors(cudaMalloc((void **)&d_a, nbytes));
// CHECK: hipStream_t *streams = (hipStream_t *) malloc(nstreams * sizeof(hipStream_t));
// allocate and initialize an array of stream handles
cudaStream_t *streams = (cudaStream_t *) malloc(nstreams * sizeof(cudaStream_t));
for (int i = 0; i < nstreams; i++)
{
// CHECK: checkCudaErrors(hipStreamCreate(&(streams[i])));
checkCudaErrors(cudaStreamCreate(&(streams[i])));
}
// CHECK: hipEvent_t start_event, stop_event;
// create CUDA event handles
cudaEvent_t start_event, stop_event;
// CHECK: checkCudaErrors(hipEventCreate(&start_event));
// CHECK: checkCudaErrors(hipEventCreate(&stop_event));
checkCudaErrors(cudaEventCreate(&start_event));
checkCudaErrors(cudaEventCreate(&stop_event));
// the events are used for synchronization only and hence do not need to record timings
// this also makes events not introduce global sync points when recorded which is critical to get overlap
// CHECK: hipEvent_t *kernelEvent;
// CHECK: kernelEvent = (hipEvent_t *) malloc(nkernels * sizeof(hipEvent_t));
cudaEvent_t *kernelEvent;
kernelEvent = (cudaEvent_t *) malloc(nkernels * sizeof(cudaEvent_t));
for (int i = 0; i < nkernels; i++)
{
// CHECK: checkCudaErrors(hipEventCreateWithFlags(&(kernelEvent[i]), hipEventDisableTiming));
checkCudaErrors(cudaEventCreateWithFlags(&(kernelEvent[i]), cudaEventDisableTiming));
}
//////////////////////////////////////////////////////////////////////
// time execution with nkernels streams
clock_t total_clocks = 0;
#if defined(__arm__) || defined(__aarch64__)
// the kernel takes more time than the channel reset time on arm archs, so to prevent hangs reduce time_clocks.
clock_t time_clocks = (clock_t)(kernel_time * (deviceProp.clockRate / 1000));
#else
clock_t time_clocks = (clock_t)(kernel_time * deviceProp.clockRate);
#endif
// CHECK: hipEventRecord(start_event, 0);
cudaEventRecord(start_event, 0);
// queue nkernels in separate streams and record when they are done
for (int i=0; i<nkernels; ++i)
{
// CHECK: hipLaunchKernelGGL(clock_block, dim3(1), dim3(1), 0, streams[i], &d_a[i], time_clocks);
clock_block<<<1,1,0,streams[i]>>>(&d_a[i], time_clocks);
total_clocks += time_clocks;
// CHECK: checkCudaErrors(hipEventRecord(kernelEvent[i], streams[i]));
checkCudaErrors(cudaEventRecord(kernelEvent[i], streams[i]));
// make the last stream wait for the kernel event to be recorded
// CHECK: checkCudaErrors(hipStreamWaitEvent(streams[nstreams-1], kernelEvent[i],0));
checkCudaErrors(cudaStreamWaitEvent(streams[nstreams-1], kernelEvent[i],0));
}
// queue a sum kernel and a copy back to host in the last stream.
// the commands in this stream get dispatched as soon as all the kernel events have been recorded
// CHECK: hipLaunchKernelGGL(sum, dim3(1), dim3(32), 0, streams[nstreams-1], d_a, nkernels);
// CHECK: checkCudaErrors(hipMemcpyAsync(a, d_a, sizeof(clock_t), hipMemcpyDeviceToHost, streams[nstreams-1]));
sum<<<1,32,0,streams[nstreams-1]>>>(d_a, nkernels);
checkCudaErrors(cudaMemcpyAsync(a, d_a, sizeof(clock_t), cudaMemcpyDeviceToHost, streams[nstreams-1]));
// at this point the CPU has dispatched all work for the GPU and can continue processing other tasks in parallel
// in this sample we just wait until the GPU is done
// CHECK: checkCudaErrors(hipEventRecord(stop_event, 0));
// CHECK: checkCudaErrors(hipEventSynchronize(stop_event));
// CHECK: checkCudaErrors(hipEventElapsedTime(&elapsed_time, start_event, stop_event));
checkCudaErrors(cudaEventRecord(stop_event, 0));
checkCudaErrors(cudaEventSynchronize(stop_event));
checkCudaErrors(cudaEventElapsedTime(&elapsed_time, start_event, stop_event));
printf("Expected time for serial execution of %d kernels = %.3fs\n", nkernels, nkernels * kernel_time/1000.0f);
printf("Expected time for concurrent execution of %d kernels = %.3fs\n", nkernels, kernel_time/1000.0f);
printf("Measured time for sample = %.3fs\n", elapsed_time/1000.0f);
bool bTestResult = (a[0] > total_clocks);
// release resources
for (int i = 0; i < nkernels; i++)
{
// CHECK: hipStreamDestroy(streams[i]);
// CHECK: hipEventDestroy(kernelEvent[i]);
cudaStreamDestroy(streams[i]);
cudaEventDestroy(kernelEvent[i]);
}
free(streams);
free(kernelEvent);
// CHECK: hipEventDestroy(start_event);
// CHECK: hipEventDestroy(stop_event);
// CHECK: hipHostFree(a);
// CHECK: hipFree(d_a);
cudaEventDestroy(start_event);
cudaEventDestroy(stop_event);
cudaFreeHost(a);
cudaFree(d_a);
if (!bTestResult)
{
printf("Test failed!\n");
exit(EXIT_FAILURE);
}
printf("Test passed\n");
exit(EXIT_SUCCESS);
}
projects/hip/tests/hipify-clang/cudaRegister.cu
-2
View File
@@ -38,7 +38,6 @@ if(status != cudaSuccess) { \
}
__global__ void Inc1(float *Ad, float *Bd){
// CHECK: int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
int tx = threadIdx.x + blockIdx.x * blockDim.x;
if(tx < 1 ){
for(int i=0;i<ITER;i++){
@@ -51,7 +50,6 @@ __global__ void Inc1(float *Ad, float *Bd){
}
__global__ void Inc2(float *Ad, float *Bd){
// CHECK: int tx = hipThreadIdx_x + hipBlockIdx_x * hipBlockDim_x;
int tx = threadIdx.x + blockIdx.x * blockDim.x;
if(tx < 1024){
for(int i=0;i<ITER;i++){
projects/hip/tests/hipify-clang/headers_test_01.cu
+6
View File
@@ -0,0 +1,6 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// CHECK: #include <hip/hip_runtime.h>
#include <cuda.h>
// CHECK-NOT: #include<cuda_runtime.h>
#include <cuda_runtime.h>
projects/hip/tests/hipify-clang/headers_test_02.cu
+6
View File
@@ -0,0 +1,6 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// CHECK: #include <hip/hip_runtime.h>
#include <cuda_runtime.h>
// CHECK-NOT: #include<cuda.h>
#include <cuda.h>
projects/hip/tests/hipify-clang/headers_test_03.cu
+10
View File
@@ -0,0 +1,10 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// CHECK: #pragma once
// CHECK-NEXT: #include <hip/hip_runtime.h>
#pragma once
// CHECK-NOT: #include <hip/hip_runtime.h>
int main(int argc, char* argv[]) {
return 0;
}
projects/hip/tests/hipify-clang/headers_test_04.cu
+12
View File
@@ -0,0 +1,12 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// CHECK: #include <hip/hip_runtime.h>
// CHECK-NEXT: #include <stdio.h>
// CHECK-NEXT: #include <iostream>
#include <stdio.h>
#include <iostream>
// CHECK-NOT: #include <hip/hip_runtime.h>
int main(int argc, char* argv[]) {
return 0;
}
projects/hip/tests/hipify-clang/headers_test_05.cu
+12
View File
@@ -0,0 +1,12 @@
// RUN: %run_test hipify "%s" "%t" %cuda_args
// CHECK: #pragma once
// CHECK-NEXT: #include <hip/hip_runtime.h>
#pragma once
// CHECK-NOT: #include <hip/hip_runtime.h>
#include <stdio.h>
int main(int argc, char* argv[]) {
return 0;
}
projects/hip/tests/hipify-clang/lit.cfg
+14 -5
View File
@@ -13,7 +13,7 @@ import lit.util
config.name = 'hipify'
# suffixes: CUDA source is only supported
config.suffixes = ['.cu']
config.suffixes = ['.cu','.cuh','.cpp','.c','.hpp','.h']
# testFormat: The test format to use to interpret tests.
config.test_format = lit.formats.ShTest()
@@ -44,8 +44,17 @@ if obj_root is not None:
path = os.path.pathsep.join((llvm_tools_dir, config.environment['PATH']))
config.environment['PATH'] = path
config.substitutions.append(("hipify", obj_root+"/hipify-clang"))
hipify_path = obj_root
clang_args = "-x cuda -v --cuda-gpu-arch=sm_30 --cuda-path='%s'"
# Clang args for CUDA...
config.substitutions.append(("%cuda_args", "-x cuda --cuda-path=%s --cuda-gpu-arch=sm_30 -isystem%s/samples/common/inc" % (config.cuda_root, config.cuda_root)))
config.substitutions.append(("%run_test", config.test_source_root + "/run_test.sh"))
if sys.platform in ['win32']:
run_test_ext = ".bat"
hipify_path += "/" + config.build_type
clang_args += " -isystem'%s'/common/inc -std=c++14"
else:
run_test_ext = ".sh"
clang_args += " -isystem'%s'/samples/common/inc"
config.substitutions.append(("%cuda_args", clang_args % (config.cuda_root, config.cuda_sdk_root)))
config.substitutions.append(("hipify", '"' + hipify_path + "/hipify-clang" + '"'))
config.substitutions.append(("%run_test", '"' + config.test_source_root + "/run_test" + run_test_ext + '"'))
projects/hip/tests/hipify-clang/lit.site.cfg.in
+15
View File
@@ -1,8 +1,23 @@
import sys
import os
config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@"
config.obj_root = "@CMAKE_CURRENT_BINARY_DIR@"
config.cuda_root = "@CUDA_TOOLKIT_ROOT_DIR@"
if sys.platform in ['win32']:
config.cuda_sdk_root = "@CUDA_SDK_ROOT_DIR@"
if not config.cuda_sdk_root or config.cuda_sdk_root == "CUDA_SDK_ROOT_DIR-NOTFOUND":
cuda_version = "@CUDA_VERSION@"
cuda_version = cuda_version.replace('.','_')
config.cuda_samples_root = os.environ.get('NVCUDASAMPLES' + cuda_version + '_ROOT')
if not config.cuda_samples_root:
lit_config.fatal('No CUDA Samples dir set! Please set CUDA_SDK_ROOT_DIR.')
config.cuda_sdk_root = config.cuda_samples_root
config.build_type = "@CMAKE_BUILD_TYPE@"
if not config.build_type:
config.build_type = "Debug"
else:
config.cuda_sdk_root = config.cuda_root
# Support substitution of the tools and libs dirs with user parameters. This is
# used when we can't determine the tool dir at configuration time.
projects/hip/tests/hipify-clang/run_test.bat
+18
View File
@@ -0,0 +1,18 @@
@echo off
setlocal
for %%i in (FileCheck.exe) do set FILE_CHECK=%%~$PATH:i
if not defined FILE_CHECK (echo Error: FileCheck.exe not found in PATH. && exit /b 1)
set HIPIFY=%1
set IN_FILE=%2
set TMP_FILE=%3
set all_args=%*
call set clang_args=%%all_args:*%4=%%
set clang_args=%4%clang_args%
%HIPIFY% -o=%TMP_FILE% %IN_FILE% -- %clang_args%
if errorlevel 1 (echo Error: hipify-clang.exe failed with exit code: %errorlevel% && exit /b %errorlevel%)
%FILE_CHECK% %IN_FILE% -input-file=%TMP_FILE%
if errorlevel 1 (echo Error: FileCheck.exe failed with exit code: %errorlevel% && exit /b %errorlevel%)
projects/hip/tests/hipify-clang/run_test.sh
-11
View File
@@ -13,16 +13,5 @@ shift 3
# Remaining args are the ones to forward to clang proper.
# Time for the classic insane little trick for making colour output work.
# A self-deleting shell-script that does the thing we want to do...
TMP_SCRIPT=$(mktemp)
cat << EOF > $TMP_SCRIPT
set -o errexit
set -o xtrace
rm $TMP_SCRIPT
$HIPIFY -o=$TMP_FILE $IN_FILE -- $@ && cat $TMP_FILE | sed -Ee 's|//.+|// |g' | FileCheck $IN_FILE
EOF
chmod a+x $TMP_SCRIPT
# Run the script via socat, spawning a virtual terminal and propagating exit code, and hence failure.
socat -du EXEC:$TMP_SCRIPT,pty,stderr STDOUT
projects/hip/tests/hipify-clang/square.cu
-2
View File
@@ -41,8 +41,6 @@ template <typename T>
__global__ void
vector_square(T *C_d, const T *A_d, size_t N)
{
// CHECK: size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x);
// CHECK: size_t stride = hipBlockDim_x * hipGridDim_x;
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x;
projects/hip/tests/src/deviceLib/hip_anyall.cpp
+1 -1
View File
@@ -21,7 +21,7 @@ THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* BUILD: %t %s ../test_common.cpp NVCC_OPTIONS --Wno-deprecated-declarations
* RUN: %t
* HIT_END
*/
projects/hip/tests/src/deviceLib/hip_ballot.cpp
+1 -1
View File
@@ -18,7 +18,7 @@ THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../test_common.cpp
* BUILD: %t %s ../test_common.cpp NVCC_OPTIONS --Wno-deprecated-declarations
* RUN: %t
* HIT_END
*/
projects/hip/tests/src/kernel/inline_asm_vadd.cpp
+1 -1
View File
@@ -16,7 +16,7 @@ IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTI
THE USE OR OTHER DEALINGS IN THE SOFTWARE. */
/* HIT_START
* BUILD: %t %s EXCLUDE_HIP_PLATFORM nvcc
* BUILD: %t %s
* RUN: %t
* HIT_END
*/
projects/hip/tests/src/p2p/hipPeerToPeer_simple.cpp
+20 -22
View File
@@ -397,32 +397,30 @@ int main(int argc, char *argv[])
if (gpuCount < 2)
{
printf("P2P application requires atleast 2 gpu devices\n");
return 0;
}
} else {
if (p_tests & 0x100) {
testPeerHostToDevice(false/*useAsyncCopy*/);
}
testPeerHostToDevice(true/*useAsyncCopy*/);
if (p_tests & 0x100) {
testPeerHostToDevice(false/*useAsyncCopy*/);
}
testPeerHostToDevice(true/*useAsyncCopy*/);
if (p_tests & 0x1) {
enablePeerFirst(false/*useAsyncCopy*/);
}
if (p_tests & 0x1) {
enablePeerFirst(false/*useAsyncCopy*/);
}
if (p_tests & 0x2) {
allocMemoryFirst(false/*useAsyncCopy*/);
}
if (p_tests & 0x2) {
allocMemoryFirst(false/*useAsyncCopy*/);
}
if (p_tests & 0x4) {
simpleNegative();
}
if (p_tests & 0x4) {
simpleNegative();
if (p_tests & 0x8) {
enablePeerFirst(true/*useAsyncCopy*/);
}
if (p_tests & 0x10) {
allocMemoryFirst(true/*useAsyncCopy*/);
}
}
if (p_tests & 0x8) {
enablePeerFirst(true/*useAsyncCopy*/);
}
if (p_tests & 0x10) {
allocMemoryFirst(true/*useAsyncCopy*/);
}
passed();
}
projects/hip/tests/src/runtimeApi/device/hipGetDeviceAttribute.cpp
+1 -1
View File
@@ -23,7 +23,7 @@ THE SOFTWARE.
/* HIT_START
* BUILD: %t %s ../../test_common.cpp
* RUN: %t EXCLUDE_HIP_PLATFORM nvcc
* RUN: %t
* HIT_END
*/
projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoD.cpp
+5 -5
View File
@@ -23,7 +23,7 @@ THE SOFTWARE.
*/
/* HIT_START
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
* BUILD: %t %s ../../test_common.cpp
* RUN: %t
* HIT_END
*/
@@ -32,7 +32,7 @@ THE SOFTWARE.
int main()
{
hipDevice_t device;
size_t Nbytes = N*sizeof(int);
int numDevices = 0;
int *A_d, *B_d, *C_d, *X_d, *Y_d, *Z_d;
@@ -69,8 +69,8 @@ int main()
HIPCHECK(hipSetDevice(1));
HIPCHECK(hipMemcpyDtoD(X_d, A_d, Nbytes));
HIPCHECK(hipMemcpyDtoD(Y_d, B_d, Nbytes));
HIPCHECK(hipMemcpyDtoD((hipDeviceptr_t)X_d, (hipDeviceptr_t)A_d, Nbytes));
HIPCHECK(hipMemcpyDtoD((hipDeviceptr_t)Y_d, (hipDeviceptr_t)B_d, Nbytes));
hipLaunchKernel(
HipTest::vectorADD,
@@ -82,7 +82,7 @@ int main()
static_cast<const int*>(Y_d),
Z_d,
N);
HIPCHECK(hipMemcpyDtoH(C_h, Z_d, Nbytes));
HIPCHECK(hipMemcpyDtoH(C_h, (hipDeviceptr_t)Z_d, Nbytes));
HIPCHECK(hipDeviceSynchronize());
HipTest::checkVectorADD(A_h, B_h, C_h, N);
projects/hip/tests/src/runtimeApi/memory/hipMemcpyDtoDAsync.cpp
+3 -4
View File
@@ -32,7 +32,6 @@ THE SOFTWARE.
int main()
{
hipDevice_t device;
size_t Nbytes = N*sizeof(int);
int numDevices = 0;
int *A_d, *B_d, *C_d, *X_d, *Y_d, *Z_d;
@@ -70,8 +69,8 @@ int main()
HIPCHECK(hipStreamCreate(&s));
HIPCHECK(hipSetDevice(1));
HIPCHECK(hipMemcpyDtoDAsync(X_d, A_d, Nbytes, s));
HIPCHECK(hipMemcpyDtoDAsync(Y_d, B_d, Nbytes, s));
HIPCHECK(hipMemcpyDtoDAsync((hipDeviceptr_t)X_d, (hipDeviceptr_t)A_d, Nbytes, s));
HIPCHECK(hipMemcpyDtoDAsync((hipDeviceptr_t)Y_d, (hipDeviceptr_t)B_d, Nbytes, s));
hipLaunchKernel(
HipTest::vectorADD,
@@ -83,7 +82,7 @@ int main()
static_cast<const int*>(Y_d),
Z_d,
N);
HIPCHECK(hipMemcpyDtoHAsync(C_h, Z_d, Nbytes, s));
HIPCHECK(hipMemcpyDtoHAsync(C_h, (hipDeviceptr_t)Z_d, Nbytes, s));
HIPCHECK(hipStreamSynchronize(s));
HIPCHECK(hipDeviceSynchronize());