Update hiphostregister test.
Move check to correct place.
[ROCm/hip commit: e7d6a34327]
Этот коммит содержится в:
@@ -51,24 +51,28 @@ int main(){
|
||||
HIPCHECK(hipHostGetDevicePointer((void**)&Ad[i], A, 0));
|
||||
}
|
||||
|
||||
// Use device pointer inside a kernel:
|
||||
// Reference the registered device pointer Ad from inside the kernel:
|
||||
for(int i=0;i<num_devices;i++){
|
||||
HIPCHECK(hipSetDevice(i));
|
||||
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/512), dim3(512), 0, 0, Ad[i]);
|
||||
hipLaunchKernel(Inc, dim3(N/512), dim3(512), 0, 0, Ad[i]);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
HIPASSERT(A[10] == 1.0f + float(num_devices));
|
||||
|
||||
|
||||
{
|
||||
// Senstizes HIP bug if device does not match where the memory was registered.
|
||||
// Sensitize HIP bug if device does not match where the memory was registered.
|
||||
HIPCHECK(hipSetDevice(0));
|
||||
|
||||
|
||||
|
||||
// Copy to B, this should be optimal pinned malloc copy:
|
||||
// Note we are using the host pointer here:
|
||||
float *Bh, *Bd;
|
||||
Bh = (float*)malloc(size);
|
||||
HIPCHECK(hipMalloc(&Bd, size));
|
||||
HIPCHECK(hipMemset(Bd, 13.0f, size));
|
||||
|
||||
for(int i=0;i<N;i++){
|
||||
A[i] = float(i);
|
||||
@@ -80,7 +84,7 @@ int main(){
|
||||
HIPCHECK(hipMemcpy(Bh, Bd, size, hipMemcpyDeviceToHost));
|
||||
|
||||
#if 0
|
||||
//TODO - disable check until we update HCC to deal with registered memory pointers.
|
||||
//TODO - disable check HCC patch for registered/locked memory usin device pointers is merged.
|
||||
for(int i=0;i<N;i++){
|
||||
if (Bh[i] != A[i]) {
|
||||
printf ("mismatch at Bh[%d]=%f, A[%d]=%f\n", i, Bh[i], i, A[i]);
|
||||
@@ -96,7 +100,6 @@ int main(){
|
||||
|
||||
|
||||
|
||||
HIPASSERT(A[10] == 1.0f + float(num_devices));
|
||||
HIPCHECK(hipHostUnregister(A));
|
||||
passed();
|
||||
}
|
||||
|
||||
Ссылка в новой задаче
Block a user