AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

Merge branch 'privatestaging' of https://github.com/AMDComputeLibraries/HIP-privatestaging into privatestaging

Telusuri Sumber 
Conflicts:
	src/hip_hcc.cpp
This commit is contained in:
Ben Sander
2016-03-19 03:22:09 -05:00
orang tua dfca4faf27 27b6a620d9
melakukan 1de63bfeea
15 mengubah file dengan 605 tambahan dan 60 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
CUDA_Runtime_API_functions_supported_by_HIP.md
+270
Melihat File
@@ -0,0 +1,270 @@
**1. Device Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaChooseDevice` | | Select compute-device which best matches criteria. |
| `cudaDeviceGetAttribute` | | Returns information about the device. |
| `cudaDeviceGetByPCIBusId` | | Returns a handle to a compute device. |
| `cudaDeviceGetCacheConfig` | `hipDeviceGetCacheConfig` | Returns the preferred cache configuration for the current device. |
| `cudaDeviceGetLimit` | | Returns resource limits. |
| `cudaDeviceGetPCIBusId` | | Returns a PCI Bus Id string for the device. |
| `cudaDeviceGetSharedMemConfig` | `hipDeviceGetSharedMemConfig` | Returns the shared memory configuration for the current device. |
| `cudaDeviceGetStreamPriorityRange` | | Returns numerical values that correspond to the least and greatest stream priorities. |
| `cudaDeviceReset` | `hipDeviceReset` | Destroy all allocations and reset all state on the current device in the current process. |
| `cudaDeviceSetCacheConfig` | `hipDeviceSetCacheConfig` | Sets the preferred cache configuration for the current device. |
| `cudaDeviceSetLimit` | | Set resource limits. |
| `cudaDeviceSetSharedMemConfig` | `hipDeviceSetSharedMemConfig` | Sets the shared memory configuration for the current device. |
| `cudaDeviceSynchronize` | `hipDeviceSynchronize` | Wait for compute device to finish. |
| `cudaGetDevice` | `hipGetDevice` | Returns which device is currently being used. |
| `cudaGetDeviceCount` | `hipGetDeviceCount` | Returns the number of compute-capable devices. |
| `cudaGetDeviceFlags` | | Gets the flags for the current device. |
| `cudaGetDeviceProperties` | `hipDeviceGetProperties` | Returns information about the compute-device. |
| `cudaIpcCloseMemHandle` | | Close memory mapped with cudaIpcOpenMemHandle. |
| `cudaIpcGetEventHandle` | | Gets an interprocess handle for a previously allocated event. |
| `cudaIpcGetMemHandle` | | Gets an interprocess memory handle for an existing device memory allocation. |
| `cudaIpcOpenEventHandle` | | Opens an interprocess event handle for use in the current process. |
| `cudaIpcOpenMemHandle` | | Opens an interprocess memory handle exported from another process and returns a device pointer usable in the local process. |
| `cudaSetDevice` | `hipSetDevice` | Set device to be used for GPU executions. |
| `cudaSetDeviceFlags` | | Sets flags to be used for device executions. |
| `cudaSetValidDevices` | | Set a list of devices that can be used for CUDA. |`
**2. Error Handling**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaGetErrorName` | `hipGetErrorName` | Returns the string representation of an error code enum name. |
| `cudaGetErrorString` | `hipGetErrorString` | Returns the description string for an error code. |
| `cudaGetLastError` | `hipGetLastError` | Returns the last error from a runtime call. |
| `cudaPeekAtLastError` | `hipPeekAtLastError` | Returns the last error from a runtime call. |
**3. Stream Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaStreamAddCallback` | | Add a callback to a compute stream. |
| `cudaStreamAttachMemAsync` | | Attach memory to a stream asynchronously. |
| `cudaStreamCreate` | `hipStreamCreate` | Create an asynchronous stream. |
| `cudaStreamCreateWithFlags` | `hipStreamCreateWithFlags` | Create an asynchronous stream. |
| `cudaStreamCreateWithPriority` | | Create an asynchronous stream with the specified priority. |
| `cudaStreamDestroy` | `hipStreamDestroy` | Destroys and cleans up an asynchronous stream. |
| `cudaStreamGetFlags` | | Query the flags of a stream. |
| `cudaStreamGetPriority` | | Query the priority of a stream. |
| `cudaStreamQuery` | | Queries an asynchronous stream for completion status. |
| `cudaStreamSynchronize` | `hipStreamSynchronize` | Waits for stream tasks to complete. |
| `cudaStreamWaitEvent` | `hipStreamWaitEvent` | Make a compute stream wait on an event. |
**4. Event Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaEventCreate` | `hipEventCreate` | Creates an event object. |
| `cudaEventCreateWithFlags` | `hipEventCreateWithFlags` | Creates an event object with the specified flags. |
| `cudaEventDestroy` | `hipEventDestroy` | Destroys an event object. |
| `cudaEventElapsedTime` | `hipEventElapsedTime` | Computes the elapsed time between events. |
| `cudaEventQuery` | `hipEventQuery` | Queries an event's status. |
| `cudaEventRecord` | `hipEventRecord` | Records an event. |
| `cudaEventSynchronize` | `hipEventSynchronize` | Waits for an event to complete. |
**5. Execution Control**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaFuncGetAttributes` | | Find out attributes for a given function. |
| `cudaFuncSetCacheConfig` | `hipFuncSetCacheConfig` | Sets the preferred cache configuration for a device function. |
| `cudaFuncSetSharedMemConfig` | | Sets the shared memory configuration for a device function. |
| `cudaGetParameterBuffer` | | Obtains a parameter buffer. |
| `cudaGetParameterBufferV2` | | Launches a specified kernel. |
| `cudaLaunchKernel` | `hipLaunchKernel` | Launches a device function. |
| `cudaSetDoubleForDevice` | | Converts a double argument to be executed on a device. |
| `cudaSetDoubleForHost` | | Converts a double argument after execution on a device. |
**6. Occupancy**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaOccupancyMaxActiveBlocksPerMultiprocessor` | | Returns occupancy for a device function. |
| `cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` | | Returns occupancy for a device function with the specified flags. |
**7. Memory Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaArrayGetInfo` | | Gets info about the specified cudaArray. |
| `cudaFree` | `hipFree` | Frees memory on the device. |
| `cudaFreeArray` | | Frees an array on the device. |
| `cudaFreeHost` | `hipHostFree` | Frees page-locked memory. |
| `cudaFreeMipmappedArray` | | Frees a mipmapped array on the device. |
| `cudaGetMipmappedArrayLevel` | | Gets a mipmap level of a CUDA mipmapped array. |
| `cudaGetSymbolAddress` | | Finds the address associated with a CUDA symbol. |
| `cudaGetSymbolSize` | | Finds the size of the object associated with a CUDA symbol. |
| `cudaHostAlloc` | `hipHostAlloc` | Allocates page-locked memory on the host. |
| `cudaHostGetDevicePointer` | `hipHostGetDevicePointer` | Passes back device pointer of mapped host memory allocated by cudaHostAlloc or registered by cudaHostRegister. |
| `cudaHostGetFlags` | `hipHostGetFlags` | Passes back flags used to allocate pinned host memory allocated by cudaHostAlloc. |
| `cudaHostRegister` | | Registers an existing host memory range for use by CUDA. |
| `cudaHostUnregister` | | Unregisters a memory range that was registered with cudaHostRegister. |
| `cudaMalloc` | `hipMalloc` | Allocate memory on the device. |
| `cudaMalloc3D` | | Allocates logical 1D, 2D, or 3D memory objects on the device. |
| `cudaMalloc3DArray` | | Allocate an array on the device. |
| `cudaMallocArray` | | Allocate an array on the device. |
| `cudaMallocHost` | `hipHostAlloc` | Allocates page-locked memory on the host. |
| `cudaMallocManaged` | | Allocates memory that will be automatically managed by the Unified Memory system. |
| `cudaMallocMipmappedArray` | | Allocate a mipmapped array on the device. |
| `cudaMallocPitch` | | Allocates pitched memory on the device. |
| `cudaMemGetInfo` | | Gets free and total device memory. |
| `cudaMemcpy` | `hipMemcpy` | Copies data between host and device. |
| `cudaMemcpy2D` | | Copies data between host and device. |
| `cudaMemcpy2DArrayToArray` | | Copies data between host and device. |
| `cudaMemcpy2DAsync` | | Copies data between host and device. |
| `cudaMemcpy2DFromArray` | | Copies data between host and device. |
| `cudaMemcpy2DFromArrayAsync` | | Copies data between host and device. |
| `cudaMemcpy2DToArray` | | Copies data between host and device. |
| `cudaMemcpy2DToArrayAsync` | | Copies data between host and device. |
| `cudaMemcpy3D` | | Copies data between 3D objects. |
| `cudaMemcpy3DAsync` | | Copies data between 3D objects. |
| `cudaMemcpy3DPeer` | | Copies memory between devices. |
| `cudaMemcpy3DPeerAsync` | | Copies memory between devices asynchronously. |
| `cudaMemcpyArrayToArray` | | Copies data between host and device. |
| `cudaMemcpyAsync` | `hipMemcpyAsync` | Copies data between host and device. |
| `cudaMemcpyFromArray` | `MemcpyFromArray` | Copies data between host and device. |
| `cudaMemcpyFromArrayAsync` | | Copies data between host and device. |
| `cudaMemcpyFromSymbol` | `hipMemcpyFromSymbol` | Copies data from the given symbol on the device. |
| `cudaMemcpyFromSymbolAsync` | | Copies data from the given symbol on the device. |
| `cudaMemcpyPeer` | `hipMemcpyPeer` | Copies memory between two devices. |
| `cudaMemcpyPeerAsync` | `hipMemcpyPeerAsync` | Copies memory between two devices asynchronously. |
| `cudaMemcpyToArray` | | Copies data between host and device. |
| `cudaMemcpyToArrayAsync` | | Copies data between host and device. |
| `cudaMemcpyToSymbol` | `hipMemcpyToSymbol` | Copies data to the given symbol on the device. |
| `cudaMemcpyToSymbolAsync` | | Copies data to the given symbol on the device. |
| `cudaMemset` | `hipMemset` | Initializes or sets device memory to a value. |
| `cudaMemset2D` | | Initializes or sets device memory to a value. |
| `cudaMemset2DAsync` | | Initializes or sets device memory to a value. |
| `cudaMemset3D` | | Initializes or sets device memory to a value. |
| `cudaMemset3DAsync` | | Initializes or sets device memory to a value. |
| `cudaMemsetAsync` | `hipMemsetAsync` | Initializes or sets device memory to a value. |
| `make\_cudaExtent` | | Returns a cudaExtent based on input parameters. |
| `make\_cudaPitchedPtr` | | Returns a cudaPitchedPtr based on input parameters. |
| `make\_cudaPos` | | Returns a cudaPos based on input parameters. |
**8. Unified Addressing**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaPointerGetAttributes` | `hipPointerGetAttributes` | Returns attributes about a specified pointer. |
**9. Peer Device Memory Access**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaDeviceCanAccessPeer` | `hipDeviceCanAccessPeer` | Queries if a device may directly access a peer device's memory. |
| `cudaDeviceDisablePeerAccess` | `hipDeviceDisablePeerAccess` | Disables direct access to memory allocations on a peer device. |
| `cudaDeviceEnablePeerAccess` | `hipDeviceEnablePeerAccess` | Enables direct access to memory allocations on a peer device. |
**10. OpenGL Interoperability**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaGLGetDevices` | | Gets the CUDA devices associated with the current OpenGL context. |
| `cudaGraphicsGLRegisterBuffer` | | Registers an OpenGL buffer object. |
| `cudaGraphicsGLRegisterImage` | | Register an OpenGL texture or renderbuffer object. |
| `cudaWGLGetDevice` | | Gets the CUDA device associated with hGpu. |
**11. Graphics Interoperability**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaGraphicsMapResources` | | Map graphics resources for access by CUDA. |
| `cudaGraphicsResourceGetMappedMipmappedArray` | | Get a mipmapped array through which to access a mapped graphics resource. |
| `cudaGraphicsResourceGetMappedPointer` | | Get a device pointer through which to access a mapped graphics resource. |
| `cudaGraphicsResourceSetMapFlags` | | Set usage flags for mapping a graphics resource. |
| `cudaGraphicsSubResourceGetMappedArray` | | Get an array through which to access a subresource of a mapped graphics resource. |
| `cudaGraphicsUnmapResources` | | Unmap graphics resources. |
| `cudaGraphicsUnregisterResource` | | Unregisters a graphics resource for access by CUDA. |
**12. Texture Reference Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaBindTexture` | | Binds a memory area to a texture. |
| `cudaBindTexture2D` | | Binds a 2D memory area to a texture. |
| `cudaBindTextureToArray` | | Binds an array to a texture. |
| `cudaBindTextureToMipmappedArray` | | Binds a mipmapped array to a texture. |
| `cudaCreateChannelDesc` | | Returns a channel descriptor using the specified format. |
| `cudaGetChannelDesc` | | Get the channel descriptor of an array. |
| `cudaGetTextureAlignmentOffset` | | Get the alignment offset of a texture. |
| `cudaGetTextureReference` | | Get the texture reference associated with a symbol. |
| `cudaUnbindTexture` | | Unbinds a texture. |
**13. Surface Reference Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaBindSurfaceToArray` | | Binds an array to a surface. |
| `cudaGetSurfaceReference` | | Get the surface reference associated with a symbol. |
**14. Texture Object Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaCreateTextureObject` | | Creates a texture object. |
| `cudaDestroyTextureObject` | | Destroys a texture object. |
| `cudaGetTextureObjectResourceDesc` | | Returns a texture object's resource descriptor. |
| `cudaGetTextureObjectResourceViewDesc` | | Returns a texture object's resource view descriptor. |
| `cudaGetTextureObjectTextureDesc` | | Returns a texture object's texture descriptor. |
**15. Surface Object Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaCreateSurfaceObject` | | Creates a surface object. |
| `cudaDestroySurfaceObject` | | Destroys a surface object. |
| `cudaGetSurfaceObjectResourceDesc` | | Returns a surface object's resource descriptor Returns the resource descriptor for the surface object specified by surfObject. |
**16. Version Management**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaDriverGetVersion` | `hipDriverGetVersion` | Returns the CUDA driver version. |
| `cudaRuntimeGetVersion` | | Returns the CUDA Runtime version. |
**17. C++ API Routines (7.0 contains, 7.5 doesn’t)**
> Will not support for HIP (probably)
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaBindSurfaceToArra`y | | Binds an array to a surface. |
| `cudaBindTexture` | | Binds a memory area to a texture. |
| `cudaBindTexture2D` | | Binds a 2D memory area to a texture. |
| `cudaBindTextureToArray` | | Binds an array to a texture. |
| `cudaBindTextureToMipmappedArray` | | Binds a mipmapped array to a texture. |
| `cudaCreateChannelDesc` | | Returns a channel descriptor using the specified format. |
| `cudaEventCreate` | | Creates an event object with the specified flags. |
| `cudaFuncGetAttributes` | | Find out attributes for a given function. |
| `cudaFuncSetCacheConfig` | | Sets the preferred cache configuration for a device function. |
| `cudaGetSymbolAddress` | | Finds the address associated with a CUDA symbol |
| `cudaGetSymbolSize` | | Finds the size of the object associated with a CUDA symbol. |
| `cudaGetTextureAlignmentOffset` | | Get the alignment offset of a texture. |
| `cudaLaunch` | | Launches a device function. |
| `cudaLaunchKernel` | | Launches a device function. |
| `cudaMallocHost` | | Allocates page-locked memory on the host |
| `cudaMallocManaged` | | Allocates memory that will be automatically managed by the Unified Memory system. |
| `cudaMemcpyFromSymbol` | | Copies data from the given symbol on the device. |
| `cudaMemcpyFromSymbolAsync` | | Copies data from the given symbol on the device. |
| `cudaMemcpyToSymbol` | | Copies data to the given symbol on the device. |
| `cudaMemcpyToSymbolAsync` | | Async copies data to the given symbol on the device. |
| `cudaOccupancyMaxActiveBlocksPerMultiprocessor` | | Returns occupancy for a device function. |
| `cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags` | | Returns occupancy for a device function with the specified flags. |
| `cudaOccupancyMaxPotentialBlockSize` | | Returns grid and block size that achieves maximum potential occupancy for a device function. |
| `cudaOccupancyMaxPotentialBlockSizeVariableSMem` | | Returns grid and block size that achieves maximum potential occupancy for a device function. |
| `cudaOccupancyMaxPotentialBlockSizeVariableSMemWithFlags` | | Returns grid and block size that achieves maximum potential occupancy for a device function. |
| `cudaOccupancyMaxPotentialBlockSizeWithFlags` | | Returns grid and block size that achived maximum potential occupancy for a device function with the specified flags. |
| `cudaSetupArgument` | | Configure a device launch. |
| `cudaStreamAttachMemAsync` | | Attach memory to a stream asynchronously. |
| `cudaUnbindTexture` | | Unbinds a texture. |
**18. Profiler Control**
| **CUDA** | **HIP** | **CUDA description** |
|-----------------------------------------------------------|-------------------------------|--------------------------------------------------------------------------------------------------------------------------------|
| `cudaProfilerInitialize` | | Initialize the CUDA profiler. |
| `cudaProfilerStart` | `hipProfilerStart` | Enable profiling. |
| `cudaProfilerStop` | `hipProfilerStop` | Disable profiling. |
samples/1_Utils/hipBusBandwidth/hipBusBandwidth.cpp
+4 -4
Melihat File
@@ -108,7 +108,7 @@ void RunBenchmark_H2D(ResultDatabase &resultDB)
float *hostMem = NULL;
if (p_pinned)
{
hipMallocHost((void**)&hostMem, sizeof(float) * numMaxFloats);
hipHostAlloc((void**)&hostMem, sizeof(float) * numMaxFloats, hipHostAllocDefault);
while (hipGetLastError() != hipSuccess)
{
// drop the size and try again
@@ -120,7 +120,7 @@ void RunBenchmark_H2D(ResultDatabase &resultDB)
return;
}
numMaxFloats = 1024 * (sizes[nSizes-1]) / 4;
hipMallocHost((void**)&hostMem, sizeof(float) * numMaxFloats);
hipHostAlloc((void**)&hostMem, sizeof(float) * numMaxFloats, hipHostAllocDefault);
}
}
else
@@ -259,9 +259,9 @@ void RunBenchmark_D2H(ResultDatabase &resultDB)
float *hostMem2;
if (p_pinned)
{
hipMallocHost((void**)&hostMem1, sizeof(float)*numMaxFloats);
hipHostAlloc((void**)&hostMem1, sizeof(float)*numMaxFloats, hipHostAllocDefault);
hipError_t err1 = hipGetLastError();
hipMallocHost((void**)&hostMem2, sizeof(float)*numMaxFloats);
hipHostAlloc((void**)&hostMem2, sizeof(float)*numMaxFloats, hipHostAllocDefault);
hipError_t err2 = hipGetLastError();
while (err1 != hipSuccess || err2 != hipSuccess)
{
samples/1_Utils/hipDispatchLatency/Makefile
+15
Melihat File
@@ -0,0 +1,15 @@
HIP_PATH?=$(shell hipconfig -p)
HIPCC=$(HIP_PATH)/bin/hipcc
EXE=hipDispatchLatency
all: install
$(EXE): hipDispatchLatency.cpp
$(HIPCC) hipDispatchLatency.cpp -o $@
install: $(EXE)
cp $(EXE) $(HIP_PATH)/bin
clean:
rm -f *.o $(EXE)
samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp
+131
Melihat File
@@ -0,0 +1,131 @@
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include"hip_runtime.h"
#include<iostream>
#include<time.h>
#define check(msg, status) \
if(status != hipSuccess){ \
printf("%s failed.\n",#msg); \
exit(1); \
}
#define LEN 1024*1024
#define SIZE LEN * sizeof(float)
#define ITER 10000
__global__ void One(hipLaunchParm lp, float* Ad){
}
int main(){
hipError_t err;
float *A, *Ad;
A = new float[LEN];
for(int i=0;i<LEN;i++){
A[i] = 1.0f;
}
hipStream_t stream;
err = hipStreamCreate(&stream);
check("Creating stream",err);
err = hipMalloc(&Ad, SIZE);
check("Allocating Ad memory on device", err);
err = hipMemcpy(Ad, A, SIZE, hipMemcpyHostToDevice);
check("Doing memory copy from A to Ad", err);
float mS = 0;
hipEvent_t start, stop;
hipEventCreate(&start);
hipEventCreate(&stop);
hipEventRecord(start);
hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
hipEventRecord(stop);
hipEventElapsedTime(&mS, start, stop);
std::cout<<"First Kernel Launch: \t\t"<<mS*1000<<" uS"<<std::endl;
hipEventRecord(start);
hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
hipEventRecord(stop);
hipEventElapsedTime(&mS, start, stop);
std::cout<<"Second Kernel Launch: \t\t"<<mS*1000<<" uS"<<std::endl;
hipEventRecord(start);
for(int i=0;i<ITER;i++){
hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
}
hipDeviceSynchronize();
hipEventRecord(stop);
hipEventElapsedTime(&mS, start, stop);
std::cout<<"NULL Stream Sync dispatch wait: \t"<<mS*1000/ITER<<" uS"<<std::endl;
hipDeviceSynchronize();
hipEventRecord(start);
for(int i=0;i<ITER;i++){
hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
}
hipEventRecord(stop);
hipDeviceSynchronize();
hipEventElapsedTime(&mS, start, stop);
std::cout<<"NULL Stream Async dispatch wait: \t"<<mS*1000/ITER<<" uS"<<std::endl;
hipDeviceSynchronize();
hipEventRecord(start);
for(int i=0;i<ITER;i++){
hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, stream, Ad);
hipDeviceSynchronize();
}
hipEventRecord(stop);
hipEventElapsedTime(&mS, start, stop);
std::cout<<"Stream Sync dispatch wait: \t\t"<<mS*1000/ITER<<" uS"<<std::endl;
hipDeviceSynchronize();
hipEventRecord(start);
for(int i=0;i<ITER;i++){
hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, stream, Ad);
}
hipDeviceSynchronize();
hipEventRecord(stop);
hipEventElapsedTime(&mS, start, stop);
std::cout<<"Stream Async dispatch wait: \t\t"<<mS*1000/ITER<<" uS"<<std::endl;
hipDeviceSynchronize();
hipEventRecord(start);
for(int i=0;i<ITER;i++){
hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, 0, Ad);
}
hipEventRecord(stop);
hipEventElapsedTime(&mS, start, stop);
std::cout<<"NULL Stream Dispatch No Wait: \t\t"<<mS*1000/ITER<<" uS"<<std::endl;
hipDeviceSynchronize();
hipEventRecord(start);
for(int i=0;i<ITER;i++){
hipLaunchKernel(HIP_KERNEL_NAME(One), dim3(LEN/512), dim3(512), 0, stream, Ad);
}
hipEventRecord(stop);
hipEventElapsedTime(&mS, start, stop);
std::cout<<"Stream Dispatch No Wait: \t\t"<<mS*1000/ITER<<" uS"<<std::endl;
hipDeviceSynchronize();
}
src/hip_hcc.cpp
+2 -7
Melihat File
@@ -2111,13 +2111,8 @@ hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags){
hip_status = hipErrorMemoryAllocation;
}else{
hc::am_memtracker_update(*ptr, device->_device_index, flags);
// void *srcPtr;
// hsa_status_t hsa_status = hsa_amd_memory_lock((*ptr), sizeBytes, &device->_hsa_agent, 1, &srcPtr);
// assert(hsa_status == HSA_STATUS_SUCCESS);
// hc::am_memtracker_add(srcPtr, sizeBytes, device->_acc, false);
}
tprintf(DB_MEM, " %s: pinned ptr=%p\n", __func__, *ptr);
}
}
return ihipLogStatus(hip_status);
}
@@ -2180,10 +2175,10 @@ hipError_t hipHostRegister(void *hostPtr, size_t sizeBytes, unsigned int flags)
return ihipLogStatus(hipErrorInvalidValue);
}
if(device){
if(flags == hipHostAllocDefault){
if(flags == hipHostRegisterDefault){
hsa_status_t hsa_status = hsa_amd_memory_lock(hostPtr, sizeBytes, &device->_hsa_agent, 1, &srcPtr);
if(hsa_status == HSA_STATUS_SUCCESS){
hip_status = hipSuccess;
hip_status = hipSuccess;
}else{
hip_status = hipErrorMemoryAllocation;
}
tests/src/cudaRegister.cu
+90
Melihat File
@@ -0,0 +1,90 @@
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include<cuda.h>
#include<cuda_runtime.h>
#include<iostream>
#include<unistd.h>
#include<stdio.h>
#include<malloc.h>
#define LEN 1024
#define SIZE LEN * sizeof(float)
#define ITER 1024*1024
#define check(msg, status){ \
if(status != cudaSuccess) { \
printf("%s failed. \n", #msg); \
} \
}
__global__ void Inc1(float *Ad, float *Bd){
int tx = threadIdx.x + blockIdx.x * blockDim.x;
if(tx < 1 ){
for(int i=0;i<ITER;i++){
Ad[tx] = Ad[tx] + 1.0f;
for(int j=0;j<256;j++){
Bd[tx] = Ad[tx];
}
}
}
}
__global__ void Inc2(float *Ad, float *Bd){
int tx = threadIdx.x + blockIdx.x * blockDim.x;
if(tx < 1024){
for(int i=0;i<ITER;i++){
Ad[tx] = Ad[tx] + 1.0f;
for(int j=0;j<256;j++){
Bd[tx] = Ad[tx];
}
}
}
}
int main(){
float *A, *Ad, *Bd;
A = new float[LEN];
for(int i=0;i<LEN;i++){
A[i] = 0.0f;
}
cudaError_t status;
status = cudaHostRegister(A, SIZE, cudaHostRegisterMapped);
check("Registering A",status);
cudaHostGetDevicePointer(&Ad, A, 0);
cudaMalloc((void**)&Bd, SIZE);
dim3 dimGrid(LEN/512,1,1);
dim3 dimBlock(512,1,1);
Inc1<<<dimGrid, dimBlock>>>(Ad, Bd);
sleep(3);
A[0] = -(ITER*1.0f);
std::cout<<"Same cache line before completion: \t"<< A[0]<<std::endl;
cudaDeviceSynchronize();
std::cout<<"Same cache line after completion: \t"<< A[0]<<std::endl;
for(int i=0;i<LEN;i++){
A[i] = 0.0f;
}
Inc2<<<dimGrid, dimBlock>>>(Ad, Bd);
sleep(3);
A[0] = -(ITER*1.0f);
std::cout<<"Diff cache line before completion: \t"<<A[0]<<std::endl;
cudaDeviceSynchronize();
std::cout<<"Diff cache line after completion: \t"<<A[0]<<std::endl;
}
tests/src/hipHostAlloc.cpp
+6 -3
Melihat File
@@ -39,14 +39,14 @@ int device;
HIPCHECK(hipGetDevice(&device));
HIPCHECK(hipGetDeviceProperties(&prop, device));
if(prop.canMapHostMemory != 1){
std::cout<<"Exiting..."<<std::endl;
//std::cout<<"Exiting..."<<std::endl;
}
HIPCHECK(hipHostAlloc((void**)&A, SIZE, hipHostAllocWriteCombined | hipHostAllocMapped));
HIPCHECK(hipHostAlloc((void**)&B, SIZE, hipHostAllocWriteCombined | hipHostAllocMapped));
HIPCHECK(hipHostAlloc((void**)&B, SIZE, hipHostAllocDefault));
HIPCHECK(hipHostAlloc((void**)&C, SIZE, hipHostAllocMapped));
HIPCHECK(hipHostGetDevicePointer((void**)&Ad, A, 0));
HIPCHECK(hipHostGetDevicePointer((void**)&Bd, B, 0));
HIPCHECK(hipHostGetDevicePointer((void**)&Cd, C, 0));
for(int i=0;i<LEN;i++){
@@ -54,6 +54,9 @@ A[i] = 1.0f;
B[i] = 2.0f;
}
HIPCHECK(hipMalloc((void**)&Bd, SIZE));
HIPCHECK(hipMemcpy(Bd, B, SIZE, hipMemcpyHostToDevice));
dim3 dimGrid(LEN/512,1,1);
dim3 dimBlock(512,1,1);
tests/src/hipMemcpy.cpp
+3 -3
Melihat File
@@ -36,7 +36,7 @@ void printSep()
// The subroutine allocates memory , copies to device, runs a vector add kernel, copies back, and checks the result.
//
// IN: numElements controls the number of elements used for allocations.
// IN: usePinnedHost : If true, allocate host with hipMallocHost and is pinned ; else allocate host memory with malloc.
// IN: usePinnedHost : If true, allocate host with hipHostAlloc and is pinned ; else allocate host memory with malloc.
// IN: useHostToHost : If true, add an extra host-to-host copy.
// IN: useDeviceToDevice : If true, add an extra deviceto-device copy after result is produced.
// IN: useMemkindDefault : If true, use memkinddefault (runtime figures out direction). if false, use explicit memcpy direction.
@@ -67,8 +67,8 @@ void memcpytest2(size_t numElements, bool usePinnedHost, bool useHostToHost, boo
if (useHostToHost) {
if (usePinnedHost) {
HIPCHECK ( hipMallocHost(&A_hh, sizeElements) );
HIPCHECK ( hipMallocHost(&B_hh, sizeElements) );
HIPCHECK ( hipHostAlloc((void**)&A_hh, sizeElements, hipHostAllocDefault) );
HIPCHECK ( hipHostAlloc((void**)&B_hh, sizeElements, hipHostAllocDefault) );
} else {
A_hh = (T*)malloc(sizeElements);
B_hh = (T*)malloc(sizeElements);
tests/src/hipMemcpyAsync.cpp
+4 -4
Melihat File
@@ -33,7 +33,7 @@ void simpleNegTest()
size_t Nbytes = N*sizeof(float);
A_malloc = (float*)malloc(Nbytes);
HIPCHECK(hipMallocHost(&A_pinned, Nbytes));
HIPCHECK(hipHostAlloc((void**)&A_pinned, Nbytes, hipHostAllocDefault));
HIPCHECK(hipMalloc(&A_d, Nbytes));
@@ -61,7 +61,7 @@ struct HostTraits<Pinned>
static void *Alloc(size_t sizeBytes) {
void *p;
HIPCHECK(hipMallocHost(&p, sizeBytes));
HIPCHECK(hipHostAlloc((void**)&p, sizeBytes, hipHostAllocDefault));
return p;
};
};
@@ -181,8 +181,8 @@ void test_manyInflightCopies(hipStream_t stream, int numElements, int numCopies,
T *A_d;
T *A_h1, *A_h2;
HIPCHECK(hipMallocHost(&A_h1, Nbytes));
HIPCHECK(hipMallocHost(&A_h2, Nbytes));
HIPCHECK(hipHostAlloc((void**)&A_h1, Nbytes, hipHostAllocDefault));
HIPCHECK(hipHostAlloc((void**)&A_h2, Nbytes, hipHostAllocDefault));
HIPCHECK(hipMalloc(&A_d, Nbytes));
for (int i=0; i<numElements; i++) {
tests/src/hipMemcpy_simple.cpp
+2 -2
Melihat File
@@ -68,8 +68,8 @@ void simpleTest2(size_t numElements, bool usePinnedHost)
T *A_d, *A_h1, *A_h2;
if (usePinnedHost) {
HIPCHECK ( hipMallocHost(&A_h1, sizeElements) );
HIPCHECK ( hipMallocHost(&A_h2, sizeElements) );
HIPCHECK ( hipHostAlloc((void**)&A_h1, sizeElements, hipHostAllocDefault) );
HIPCHECK ( hipHostAlloc((void**)&A_h2, sizeElements, hipHostAllocDefault) );
} else {
A_h1 = (T*)aligned_alloc(alignment, sizeElements);
HIPASSERT(A_h1);
tests/src/hipMultiThreadStreams2.cpp
+27 -27
Melihat File
@@ -34,21 +34,21 @@ Array[tx] = Array[tx] + T(1);
void run1(size_t size, hipStream_t stream){
float *Ah, *Bh, *Cd, *Dd, *Eh;
hipMallocHost(&Ah, size);
hipMallocHost(&Bh, size);
hipMalloc(&Cd, size);
hipMalloc(&Dd, size);
hipMallocHost(&Eh, size);
HIPCHECK(hipHostAlloc((void**)&Ah, size, hipHostAllocDefault));
HIPCHECK(hipHostAlloc((void**)&Bh, size, hipHostAllocDefault));
HIPCHECK(hipMalloc(&Cd, size));
HIPCHECK(hipMalloc(&Dd, size));
HIPCHECK(hipHostAlloc((void**)&Eh, size, hipHostAllocDefault));
for(int i=0;i<N;i++){
Ah[i] = 1.0f;
}
hipMemcpyAsync(Bh, Ah, size, hipMemcpyHostToHost, stream);
hipMemcpyAsync(Cd, Bh, size, hipMemcpyHostToDevice, stream);
HIPCHECK(hipMemcpyAsync(Bh, Ah, size, hipMemcpyHostToHost, stream));
HIPCHECK(hipMemcpyAsync(Cd, Bh, size, hipMemcpyHostToDevice, stream));
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream, Cd);
hipMemcpyAsync(Dd, Cd, size, hipMemcpyDeviceToDevice, stream);
hipMemcpyAsync(Eh, Dd, size, hipMemcpyDeviceToHost, stream);
HIPCHECK(hipMemcpyAsync(Dd, Cd, size, hipMemcpyDeviceToDevice, stream));
HIPCHECK(hipMemcpyAsync(Eh, Dd, size, hipMemcpyDeviceToHost, stream));
HIPCHECK(hipDeviceSynchronize());
HIPASSERT(Eh[10] == Ah[10] + 1.0f);
}
@@ -58,27 +58,27 @@ void run(size_t size, hipStream_t stream1, hipStream_t stream2){
float *Ah, *Bh, *Cd, *Dd, *Eh;
float *Ahh, *Bhh, *Cdd, *Ddd, *Ehh;
hipMallocHost(&Ah, size);
hipMallocHost(&Bh, size);
hipMalloc(&Cd, size);
hipMalloc(&Dd, size);
hipMallocHost(&Eh, size);
hipMallocHost(&Ahh, size);
hipMallocHost(&Bhh, size);
hipMalloc(&Cdd, size);
hipMalloc(&Ddd, size);
hipMallocHost(&Ehh, size);
HIPCHECK(hipHostAlloc((void**)&Ah, size, hipHostAllocDefault));
HIPCHECK(hipHostAlloc((void**)&Bh, size, hipHostAllocDefault));
HIPCHECK(hipMalloc(&Cd, size));
HIPCHECK(hipMalloc(&Dd, size));
HIPCHECK(hipHostAlloc((void**)&Eh, size, hipHostAllocDefault));
HIPCHECK(hipHostAlloc((void**)&Ahh, size, hipHostAllocDefault));
HIPCHECK(hipHostAlloc((void**)&Bhh, size, hipHostAllocDefault));
HIPCHECK(hipMalloc(&Cdd, size));
HIPCHECK(hipMalloc(&Ddd, size));
HIPCHECK(hipHostAlloc((void**)&Ehh, size, hipHostAllocDefault));
hipMemcpyAsync(Bh, Ah, size, hipMemcpyHostToHost, stream1);
hipMemcpyAsync(Bhh, Ahh, size, hipMemcpyHostToHost, stream2);
hipMemcpyAsync(Cd, Bh, size, hipMemcpyHostToDevice, stream1);
hipMemcpyAsync(Cdd, Bhh, size, hipMemcpyHostToDevice, stream2);
HIPCHECK(hipMemcpyAsync(Bh, Ah, size, hipMemcpyHostToHost, stream1));
HIPCHECK(hipMemcpyAsync(Bhh, Ahh, size, hipMemcpyHostToHost, stream2));
HIPCHECK(hipMemcpyAsync(Cd, Bh, size, hipMemcpyHostToDevice, stream1));
HIPCHECK(hipMemcpyAsync(Cdd, Bhh, size, hipMemcpyHostToDevice, stream2));
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream1, Cd);
hipLaunchKernel(HIP_KERNEL_NAME(Inc), dim3(N/500), dim3(500), 0, stream2, Cdd);
hipMemcpyAsync(Dd, Cd, size, hipMemcpyDeviceToDevice, stream1);
hipMemcpyAsync(Ddd, Cdd, size, hipMemcpyDeviceToDevice, stream2);
hipMemcpyAsync(Eh, Dd, size, hipMemcpyDeviceToHost, stream1);
hipMemcpyAsync(Ehh, Ddd, size, hipMemcpyDeviceToHost, stream2);
HIPCHECK(hipMemcpyAsync(Dd, Cd, size, hipMemcpyDeviceToDevice, stream1));
HIPCHECK(hipMemcpyAsync(Ddd, Cdd, size, hipMemcpyDeviceToDevice, stream2));
HIPCHECK(hipMemcpyAsync(Eh, Dd, size, hipMemcpyDeviceToHost, stream1));
HIPCHECK(hipMemcpyAsync(Ehh, Ddd, size, hipMemcpyDeviceToHost, stream2));
HIPCHECK(hipDeviceSynchronize());
HIPASSERT(Eh[10] = Ah[10] + 1.0f);
HIPASSERT(Ehh[10] = Ahh[10] + 1.0f);
tests/src/hipPerfMemcpy.cpp
+41
Melihat File
@@ -0,0 +1,41 @@
#include"test_common.h"
#include<iostream>
#include<time.h>
#define NUM_SIZE 8
#define NUM_ITER 12
static size_t size[NUM_SIZE];
void setup(){
for(int i=0;i<NUM_SIZE;i++){
size[i] = 1<<(i+6); // start at 8 bytes
}
}
void valSet(int *A, int val, size_t size){
size_t len = size/sizeof(int);
for(int i=0;i<len;i++){
A[i] = val;
}
}
int main(){
setup();
int *A, *Ad;
for(int i=0;i<NUM_SIZE;i++){
std::cout<<size[i]<<std::endl;
A = (int*)malloc(size[i]);
valSet(A, 1, size[i]);
hipMalloc(&Ad, size[i]);
std::cout<<"Malloc success at size: "<<size[i]<<std::endl;
clock_t start ,end;
start = clock();
for(int i=0;i<NUM_ITER;i++){
hipMemcpy(Ad, A, size[i], hipMemcpyHostToDevice);
}
hipDeviceSynchronize();
end = clock();
double uS = (double)(end - start)*1000/(NUM_ITER*CLOCKS_PER_SEC);
std::cout<<uS<<std::endl;
}
}
tests/src/hipPointerAttrib.cpp
+4 -4
Melihat File
@@ -115,7 +115,7 @@ void testSimple()
hipError_t e;
HIPCHECK ( hipMalloc(&A_d, Nbytes) );
HIPCHECK ( hipMallocHost(&A_Pinned_h, Nbytes) );
HIPCHECK ( hipHostAlloc((void**)&A_Pinned_h, Nbytes, hipHostAllocDefault) );
A_OSAlloc_h = (char*)malloc(Nbytes);
size_t free, total;
@@ -168,7 +168,7 @@ void testSimple()
// Device-visible host memory
printf ("\nDevice-visible host memory (hipMallocHost)\n");
printf ("\nDevice-visible host memory (hipHostAlloc)\n");
HIPCHECK( hipPointerGetAttributes(&attribs, A_Pinned_h));
printf("getAttr:%-20s", "A_pinned_h"); printAttribs(&attribs);
@@ -277,13 +277,13 @@ void clusterAllocs(int numAllocs, size_t minSize, size_t maxSize)
void * ptr;
if (isDevice) {
totalDeviceAllocated[reference[i]._attrib.device] += reference[i]._sizeBytes;
HIPCHECK(hipMalloc(&ptr, reference[i]._sizeBytes));
HIPCHECK(hipHostAlloc((void**)&ptr, reference[i]._sizeBytes, hipHostAllocDefault));
reference[i]._attrib.memoryType = hipMemoryTypeDevice;
reference[i]._attrib.devicePointer = ptr;
reference[i]._attrib.hostPointer = NULL;
reference[i]._attrib.allocationFlags = 0; // TODO-randomize these.
} else {
HIPCHECK(hipMallocHost(&ptr, reference[i]._sizeBytes));
HIPCHECK(hipHostAlloc((void**)&ptr, reference[i]._sizeBytes, hipHostAllocDefault));
reference[i]._attrib.memoryType = hipMemoryTypeHost;
reference[i]._attrib.devicePointer = ptr;
reference[i]._attrib.hostPointer = ptr;
tests/src/hipStream.h
+2 -2
Melihat File
@@ -85,7 +85,7 @@ void initArrays(T **Ad, T **Ah,
HIPCHECK( hipMalloc(Ad, NBytes));
}
if(usePinnedHost){
HIPCHECK( hipMallocHost(Ah, NBytes));
HIPCHECK( hipHostAlloc((void**)Ah, NBytes, hipHostAllocDefault));
}
else{
*Ah = new T[N];
@@ -102,7 +102,7 @@ void initArrays(T **Ad, size_t N,
HIPCHECK( hipMalloc(Ad, NBytes));
}else{
if(usePinnedHost){
HIPCHECK(hipMallocHost(Ad, NBytes));
HIPCHECK(hipHostAlloc((void**)Ad, NBytes, hipHostAllocDefault));
}else{
*Ad = new T[N];
HIPASSERT(*Ad != NULL);
tests/src/test_common.h
+4 -4
Melihat File
@@ -141,13 +141,13 @@ void initArrays(T **A_d, T **B_d, T **C_d,
if (usePinnedHost) {
if (A_h) {
HIPCHECK ( hipMallocHost(A_h, Nbytes) );
HIPCHECK ( hipHostAlloc((void**)A_h, Nbytes, hipHostAllocDefault) );
}
if (B_h) {
HIPCHECK ( hipMallocHost(B_h, Nbytes) );
HIPCHECK ( hipHostAlloc((void**)B_h, Nbytes, hipHostAllocDefault) );
}
if (C_h) {
HIPCHECK ( hipMallocHost(C_h, Nbytes) );
HIPCHECK ( hipHostAlloc((void**)C_h, Nbytes, hipHostAllocDefault) );
}
} else {
if (A_h) {
@@ -258,7 +258,7 @@ struct Pinned {
static void *Alloc(size_t sizeBytes)
{
void *p;
HIPCHECK(hipMallocHost(&p, sizeBytes));
HIPCHECK(hipHostAlloc((void**)&p, sizeBytes, hipHostAllocDefault));
return p;
};
};