diff --git a/projects/clr/hipamd/bin/hipify-perl b/projects/clr/hipamd/bin/hipify-perl index 4b6e6e89b9..a37f728bd1 100755 --- a/projects/clr/hipamd/bin/hipify-perl +++ b/projects/clr/hipamd/bin/hipify-perl @@ -247,6 +247,7 @@ while (@ARGV) { $ft{'err'} += s/\bcudaErrorInvalidValue\b/hipErrorInvalidValue/g; $ft{'err'} += s/\bcudaErrorInvalidResourceHandle\b/hipErrorInvalidResourceHandle/g; $ft{'err'} += s/\bcudaErrorInvalidDevice\b/hipErrorInvalidDevice/g; + $ft{'err'} += s/\bcudaErrorInvalidDevicePointer\b/hipErrorInvalidDevicePointer/g; $ft{'err'} += s/\bcudaErrorNoDevice\b/hipErrorNoDevice/g; $ft{'err'} += s/\bcudaErrorNotReady\b/hipErrorNotReady/g; $ft{'err'} += s/\bcudaErrorUnknown\b/hipErrorUnknown/g; @@ -280,6 +281,7 @@ while (@ARGV) { $ft{'mem'} += s/\bcudaMemcpyKind\b/hipMemcpyKind/g; $ft{'mem'} += s/\bcudaPointerAttributes\b/hipPointerAttribute_t/g; + $ft{'mem'} += s/\bcudaPointerGetAttributes\b/hipPointerGetAttribute_t/g; $ft{'mem'} += s/\bcudaMemcpy2D\b/hipMemcpy2D/g; $ft{'mem'} += s/\bcudaMemcpy2DToArray\b/hipMemcpy2DToArray/g; @@ -339,6 +341,7 @@ while (@ARGV) { $ft{'event'} += s/\bcudaEventElapsedTime\b/hipEventElapsedTime/g; $ft{'event'} += s/\bcudaEventSynchronize\b/hipEventSynchronize/g; $ft{'event'} += s/\bcudaEventDisableTiming\b/hipEventDisableTiming/g; + $ft{'event'} += s/\bcudaEventQuery\b/hipEventQuery/g; #-------- # Streams @@ -489,15 +492,15 @@ while (@ARGV) { my $kernelName; # Handle the <>> syntax: - $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernel(HIP_KERNEL_NAME($1$2), dim3($3), dim3($4), $5, $6, /g; + $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), $5, $6, /g; $kernelName = $1 if $k; # Handle the <>> syntax: - $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernel(HIP_KERNEL_NAME($1$2), dim3($3), dim3($4), $5, 0, /g; + $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), $5, 0, /g; $kernelName = $1 if $k; # Handle the <>> syntax: - $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>([\s\\]*)\(/hipLaunchKernel(HIP_KERNEL_NAME($1$2), dim3($3), dim3($4), 0, 0, /g; + $k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>([\s\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), 0, 0, /g; $kernelName = $1 if $k; $ft{'kern'} += $k; diff --git a/projects/clr/hipamd/include/hip/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hip_runtime_api.h index 0cdace0e99..97733964bf 100644 --- a/projects/clr/hipamd/include/hip/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hip_runtime_api.h @@ -161,6 +161,7 @@ typedef enum hipError_t { hipErrorProfilerNotInitialized = 6, hipErrorProfilerAlreadyStarted = 7, hipErrorProfilerAlreadyStopped = 8, + hipErrorInsufficientDriver = 35, hipErrorInvalidImage = 200, hipErrorInvalidContext = 201, ///< Produced when input context is invalid. hipErrorContextAlreadyCurrent = 202, diff --git a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h index cd8b2c17a0..1103bb2274 100644 --- a/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime_api.h @@ -31,6 +31,13 @@ THE SOFTWARE. extern "C" { #endif +#ifdef __cplusplus + #define __dparm(x) \ + = x +#else + #define __dparm(x) +#endif + //TODO -move to include/hip_runtime_api.h as a common implementation. /** * Memory copy types @@ -102,8 +109,8 @@ typedef cudaEvent_t hipEvent_t; typedef cudaStream_t hipStream_t; typedef cudaIpcEventHandle_t hipIpcEventHandle_t; typedef cudaIpcMemHandle_t hipIpcMemHandle_t; -typedef cudaLimit hipLimit_t; -typedef cudaFuncCache hipFuncCache_t; +typedef enum cudaLimit hipLimit_t; +typedef enum cudaFuncCache hipFuncCache_t; typedef CUcontext hipCtx_t; typedef cudaSharedMemConfig hipSharedMemConfig; typedef CUfunc_cache hipFuncCache; @@ -112,10 +119,10 @@ typedef CUdevice hipDevice_t; typedef CUmodule hipModule_t; typedef CUfunction hipFunction_t; typedef CUdeviceptr hipDeviceptr_t; -typedef cudaChannelFormatKind hipChannelFormatKind; -typedef cudaChannelFormatDesc hipChannelFormatDesc; -typedef cudaTextureReadMode hipTextureReadMode; -typedef cudaArray hipArray; +typedef enum cudaChannelFormatKind hipChannelFormatKind; +typedef struct cudaChannelFormatDesc hipChannelFormatDesc; +typedef enum cudaTextureReadMode hipTextureReadMode; +typedef struct cudaArray hipArray; // Flags that can be used with hipStreamCreateWithFlags #define hipStreamDefault cudaStreamDefault @@ -192,7 +199,7 @@ switch(hError) { } } -inline static cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind kind) { +inline static enum cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind kind) { switch(kind) { case hipMemcpyHostToHost: return cudaMemcpyHostToHost; @@ -255,7 +262,7 @@ inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int fla return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags)); } -inline static hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags) { +inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags) { return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags)); } @@ -294,8 +301,8 @@ inline static hipError_t hipSetDevice(int device) { inline static hipError_t hipChooseDevice( int* device, const hipDeviceProp_t* prop ) { - cudaDeviceProp cdprop; - memset(&cdprop,0x0,sizeof(cudaDeviceProp)); + struct cudaDeviceProp cdprop; + memset(&cdprop,0x0,sizeof(struct cudaDeviceProp)); cdprop.major= prop->major; cdprop.minor = prop->minor; cdprop.totalGlobalMem = prop->totalGlobalMem ; @@ -356,25 +363,24 @@ inline static hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, } -inline static hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind copyKind, hipStream_t stream=0) { +inline static hipError_t hipMemcpyAsync(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind copyKind, hipStream_t stream __dparm(0)) { return hipCUDAErrorTohipError(cudaMemcpyAsync(dst, src, sizeBytes, hipMemcpyKindToCudaMemcpyKind(copyKind), stream)); } - -inline static hipError_t hipMemcpyToSymbol(const void* symbol, const void* src, size_t sizeBytes, size_t offset = 0, hipMemcpyKind copyType = hipMemcpyHostToDevice) { +inline static hipError_t hipMemcpyToSymbol(const void* symbol, const void* src, size_t sizeBytes, size_t offset __dparm(0), hipMemcpyKind copyType __dparm(hipMemcpyHostToDevice)) { return hipCUDAErrorTohipError(cudaMemcpyToSymbol(symbol, src, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(copyType))); } -inline static hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t sizeBytes, size_t offset, hipMemcpyKind copyType, hipStream_t stream = 0) { +inline static hipError_t hipMemcpyToSymbolAsync(const void* symbol, const void* src, size_t sizeBytes, size_t offset, hipMemcpyKind copyType, hipStream_t stream __dparm(0)) { return hipCUDAErrorTohipError(cudaMemcpyToSymbolAsync(symbol, src, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(copyType), stream)); } -inline static hipError_t hipMemcpyFromSymbol(void *dst, const void* symbolName, size_t sizeBytes, size_t offset = 0, hipMemcpyKind kind = hipMemcpyDeviceToHost) +inline static hipError_t hipMemcpyFromSymbol(void *dst, const void* symbolName, size_t sizeBytes, size_t offset __dparm(0), hipMemcpyKind kind __dparm(hipMemcpyDeviceToHost)) { return hipCUDAErrorTohipError(cudaMemcpyFromSymbol(dst, symbolName, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(kind))); } -inline static hipError_t hipMemcpyFromSymbolAsync(void *dst, const void* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream = 0) +inline static hipError_t hipMemcpyFromSymbolAsync(void *dst, const void* symbolName, size_t sizeBytes, size_t offset, hipMemcpyKind kind, hipStream_t stream __dparm(0)) { return hipCUDAErrorTohipError(cudaMemcpyFromSymbolAsync(dst, symbolName, sizeBytes, offset, hipMemcpyKindToCudaMemcpyKind(kind), stream)); } @@ -443,7 +449,7 @@ inline static hipError_t hipMemset(void* devPtr,int value, size_t count) { return hipCUDAErrorTohipError(cudaMemset(devPtr, value, count)); } -inline static hipError_t hipMemsetAsync(void* devPtr,int value, size_t count, hipStream_t stream = 0) { +inline static hipError_t hipMemsetAsync(void* devPtr,int value, size_t count, hipStream_t stream __dparm(0)) { return hipCUDAErrorTohipError(cudaMemsetAsync(devPtr, value, count, stream)); } @@ -454,7 +460,7 @@ inline static hipError_t hipMemsetD8(hipDeviceptr_t dest, unsigned char value, inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t *p_prop, int device) { - cudaDeviceProp cdprop; + struct cudaDeviceProp cdprop; cudaError_t cerror; cerror = cudaGetDeviceProperties(&cdprop,device); strncpy(p_prop->name,cdprop.name, 256); @@ -515,7 +521,7 @@ inline static hipError_t hipGetDeviceProperties(hipDeviceProp_t *p_prop, int dev inline static hipError_t hipDeviceGetAttribute(int* pi, hipDeviceAttribute_t attr, int device) { - cudaDeviceAttr cdattr; + enum cudaDeviceAttr cdattr; cudaError_t cerror; switch (attr) { @@ -591,7 +597,7 @@ inline static hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor( } inline static hipError_t hipPointerGetAttributes(hipPointerAttribute_t *attributes, void* ptr){ - cudaPointerAttributes cPA; + struct cudaPointerAttributes cPA; hipError_t err = hipCUDAErrorTohipError(cudaPointerGetAttributes(&cPA, ptr)); if(err == hipSuccess){ switch (cPA.memoryType){ @@ -622,7 +628,7 @@ inline static hipError_t hipEventCreate( hipEvent_t* event) return hipCUDAErrorTohipError(cudaEventCreate(event)); } -inline static hipError_t hipEventRecord( hipEvent_t event, hipStream_t stream = NULL) +inline static hipError_t hipEventRecord( hipEvent_t event, hipStream_t stream __dparm(NULL)) { return hipCUDAErrorTohipError(cudaEventRecord(event,stream)); } @@ -755,18 +761,18 @@ inline static hipError_t hipMemcpyPeer ( void* dst, int dstDevice, const void* return hipCUDAErrorTohipError(cudaMemcpyPeer(dst, dstDevice, src, srcDevice, count)); } -inline static hipError_t hipMemcpyPeerAsync ( void* dst, int dstDevice, const void* src, int srcDevice, size_t count, hipStream_t stream=0 ) +inline static hipError_t hipMemcpyPeerAsync ( void* dst, int dstDevice, const void* src, int srcDevice, size_t count, hipStream_t stream __dparm(0)) { return hipCUDAErrorTohipError(cudaMemcpyPeerAsync(dst, dstDevice, src, srcDevice, count, stream)); } // Profile APIs: -inline hipError_t hipProfilerStart() +inline static hipError_t hipProfilerStart() { return hipCUDAErrorTohipError(cudaProfilerStart()); } -inline hipError_t hipProfilerStop() +inline static hipError_t hipProfilerStop() { return hipCUDAErrorTohipError(cudaProfilerStop()); }