diff --git a/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md b/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md index 7c2d2c6631..010cf4d2c9 100644 --- a/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md +++ b/projects/clr/hipamd/docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md @@ -4,8 +4,8 @@ | **type** | **CUDA** | **HIP** | |-------------:|--------------------------------------------------------------------|------------------------------------------------------------| -| struct |`CUDA_ARRAY3D_DESCRIPTOR` | | -| typedef |`CUDA_ARRAY3D_DESCRIPTOR_st` | | +| struct |`CUDA_ARRAY3D_DESCRIPTOR` |`HIP_ARRAY3D_DESCRIPTOR` | +| typedef |`CUDA_ARRAY3D_DESCRIPTOR_st` |`HIP_ARRAY3D_DESCRIPTOR` | | struct |`CUDA_ARRAY_DESCRIPTOR` |`HIP_ARRAY_DESCRIPTOR` | | typedef |`CUDA_ARRAY_DESCRIPTOR_st` |`HIP_ARRAY_DESCRIPTOR` | | struct |`CUDA_MEMCPY2D` |`hip_Memcpy2D` | @@ -30,12 +30,12 @@ | typedef |`CUipcMemHandle_st` |`hipIpcMemHandle_st` | | union |`CUstreamBatchMemOpParams` | | | typedef |`CUstreamBatchMemOpParams_union` | | -| enum |***`CUaddress_mode`*** | | -| typedef |***`CUaddress_mode_enum`*** | | -| 0 |*`CU_TR_ADDRESS_MODE_WRAP`* | | -| 1 |*`CU_TR_ADDRESS_MODE_CLAMP`* | | -| 2 |*`CU_TR_ADDRESS_MODE_MIRROR`* | | -| 3 |*`CU_TR_ADDRESS_MODE_BORDER`* | | +| enum |***`CUaddress_mode`*** |***`hipTextureAddressMode`*** | +| typedef |***`CUaddress_mode_enum`*** |***`hipTextureAddressMode`*** | +| 0 |*`CU_TR_ADDRESS_MODE_WRAP`* |*`hipAddressModeWrap`* | +| 1 |*`CU_TR_ADDRESS_MODE_CLAMP`* |*`hipAddressModeClamp`* | +| 2 |*`CU_TR_ADDRESS_MODE_MIRROR`* |*`hipAddressModeMirror`* | +| 3 |*`CU_TR_ADDRESS_MODE_BORDER`* |*`hipAddressModeBorder`* | | enum |***`CUarray_cubemap_face`*** | | | typedef |***`CUarray_cubemap_face_enum`*** | | | 0x00 |*`CU_CUBEMAP_FACE_POSITIVE_X`* | | @@ -557,9 +557,9 @@ | define |`CU_PARAM_TR_DEFAULT` | | | define |`CU_STREAM_LEGACY` | | | define |`CU_STREAM_PER_THREAD` | | -| define |`CU_TRSA_OVERRIDE_FORMAT` | | -| define |`CU_TRSF_NORMALIZED_COORDINATES` | | -| define |`CU_TRSF_READ_AS_INTEGER` | | +| define |`CU_TRSA_OVERRIDE_FORMAT` |`HIP_TRSA_OVERRIDE_FORMAT` | +| define |`CU_TRSF_NORMALIZED_COORDINATES` |`HIP_TRSF_NORMALIZED_COORDINATES` | +| define |`CU_TRSF_READ_AS_INTEGER` |`HIP_TRSF_READ_AS_INTEGER` | | define |`CU_TRSF_SRGB` | | | define |`CUDA_ARRAY3D_2DARRAY` | | | define |`CUDA_ARRAY3D_CUBEMAP` |`hipArrayCubemap` | diff --git a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp index 232cd16130..51dc9665c3 100644 --- a/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp +++ b/projects/clr/hipamd/hipify-clang/src/CUDA2HIP_Driver_API_types.cpp @@ -27,8 +27,8 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ // 1. Structs - {"CUDA_ARRAY3D_DESCRIPTOR_st", {"HIP_ARRAY3D_DESCRIPTOR", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, - {"CUDA_ARRAY3D_DESCRIPTOR", {"HIP_ARRAY3D_DESCRIPTOR", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + {"CUDA_ARRAY3D_DESCRIPTOR_st", {"HIP_ARRAY3D_DESCRIPTOR", "", CONV_TYPE, API_DRIVER}}, + {"CUDA_ARRAY3D_DESCRIPTOR", {"HIP_ARRAY3D_DESCRIPTOR", "", CONV_TYPE, API_DRIVER}}, {"CUDA_ARRAY_DESCRIPTOR_st", {"HIP_ARRAY_DESCRIPTOR", "", CONV_TYPE, API_DRIVER}}, {"CUDA_ARRAY_DESCRIPTOR", {"HIP_ARRAY_DESCRIPTOR", "", CONV_TYPE, API_DRIVER}}, @@ -204,13 +204,13 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ {"CUstreamBatchMemOpParams_union", {"hipStreamBatchMemOpParams", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, // 3. Enums - {"CUaddress_mode", {"hipAddress_mode", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, - {"CUaddress_mode_enum", {"hipAddress_mode", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, + {"CUaddress_mode", {"hipTextureAddressMode", "", CONV_TYPE, API_DRIVER}}, + {"CUaddress_mode_enum", {"hipTextureAddressMode", "", CONV_TYPE, API_DRIVER}}, // CUaddress_mode enum values - {"CU_TR_ADDRESS_MODE_WRAP", {"HIP_TR_ADDRESS_MODE_WRAP", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 0 - {"CU_TR_ADDRESS_MODE_CLAMP", {"HIP_TR_ADDRESS_MODE_CLAMP", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 1 - {"CU_TR_ADDRESS_MODE_MIRROR", {"HIP_TR_ADDRESS_MODE_MIRROR", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 2 - {"CU_TR_ADDRESS_MODE_BORDER", {"HIP_TR_ADDRESS_MODE_BORDER", "", CONV_NUMERIC_LITERAL, API_DRIVER, HIP_UNSUPPORTED}}, // 3 + {"CU_TR_ADDRESS_MODE_WRAP", {"hipAddressModeWrap", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 0 + {"CU_TR_ADDRESS_MODE_CLAMP", {"hipAddressModeClamp", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 1 + {"CU_TR_ADDRESS_MODE_MIRROR", {"hipAddressModeMirror", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 2 + {"CU_TR_ADDRESS_MODE_BORDER", {"hipAddressModeBorder", "", CONV_NUMERIC_LITERAL, API_DRIVER}}, // 3 {"CUarray_cubemap_face", {"hipGraphicsCubeFace", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, {"CUarray_cubemap_face_enum", {"hipGraphicsCubeFace", "", CONV_TYPE, API_DRIVER, HIP_UNSUPPORTED}}, @@ -1488,9 +1488,9 @@ const std::map CUDA_DRIVER_TYPE_NAME_MAP{ {"CU_STREAM_LEGACY", {"hipStreamLegacy", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // ((CUstream)0x1) // cudaStreamPerThread ((cudaStream_t)0x2) {"CU_STREAM_PER_THREAD", {"hipStreamPerThread", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // ((CUstream)0x2) - {"CU_TRSA_OVERRIDE_FORMAT", {"HIP_TRSA_OVERRIDE_FORMAT", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01 - {"CU_TRSF_NORMALIZED_COORDINATES", {"HIP_TRSF_NORMALIZED_COORDINATES", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x02 - {"CU_TRSF_READ_AS_INTEGER", {"HIP_TRSF_READ_AS_INTEGER", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x01 + {"CU_TRSA_OVERRIDE_FORMAT", {"HIP_TRSA_OVERRIDE_FORMAT", "", CONV_DEFINE, API_DRIVER}}, // 0x01 + {"CU_TRSF_NORMALIZED_COORDINATES", {"HIP_TRSF_NORMALIZED_COORDINATES", "", CONV_DEFINE, API_DRIVER}}, // 0x02 + {"CU_TRSF_READ_AS_INTEGER", {"HIP_TRSF_READ_AS_INTEGER", "", CONV_DEFINE, API_DRIVER}}, // 0x01 {"CU_TRSF_SRGB", {"HIP_TRSF_SRGB", "", CONV_DEFINE, API_DRIVER, HIP_UNSUPPORTED}}, // 0x10 // no analogue // NOTE: Deprecated, use CUDA_ARRAY3D_LAYERED diff --git a/projects/clr/hipamd/include/hip/hcc_detail/driver_types.h b/projects/clr/hipamd/include/hip/hcc_detail/driver_types.h index 5b2297114f..18f344d642 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/driver_types.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/driver_types.h @@ -59,14 +59,21 @@ typedef enum hipArray_Format { }hipArray_Format; typedef struct HIP_ARRAY_DESCRIPTOR { - enum hipArray_Format format; - unsigned int numChannels; - size_t width; - size_t height; - unsigned int flags; - size_t depth; + size_t Width; + size_t Height; + enum hipArray_Format Format; + unsigned int NumChannels; }HIP_ARRAY_DESCRIPTOR; +typedef struct HIP_ARRAY3D_DESCRIPTOR { + size_t Width; + size_t Height; + size_t Depth; + enum hipArray_Format Format; + unsigned int NumChannels; + unsigned int Flags; +}HIP_ARRAY3D_DESCRIPTOR; + typedef struct hipArray { void* data; // FIXME: generalize this struct hipChannelFormatDesc desc; @@ -74,7 +81,8 @@ typedef struct hipArray { unsigned int width; unsigned int height; unsigned int depth; - struct HIP_ARRAY_DESCRIPTOR drvDesc; + enum hipArray_Format Format; + unsigned int NumChannels; bool isDrv; unsigned int textureType; }hipArray; diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 97df1f4219..979efebe4d 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -1821,7 +1821,7 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, si size_t height __dparm(0), unsigned int flags __dparm(hipArrayDefault)); hipError_t hipArrayCreate(hipArray** pHandle, const HIP_ARRAY_DESCRIPTOR* pAllocateArray); -hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray); +hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAllocateArray); hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent); diff --git a/projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp index 36d37a4fad..2cb9877cac 100644 --- a/projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp +++ b/projects/clr/hipamd/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp @@ -58,10 +58,10 @@ bool runTest(int argc, char** argv) { hipArray* array; HIP_ARRAY_DESCRIPTOR desc; - desc.format = HIP_AD_FORMAT_FLOAT; - desc.numChannels = 1; - desc.width = width; - desc.height = height; + desc.Format = HIP_AD_FORMAT_FLOAT; + desc.NumChannels = 1; + desc.Width = width; + desc.Height = height; hipArrayCreate(&array, &desc); hip_Memcpy2D copyParam; diff --git a/projects/clr/hipamd/src/hip_memory.cpp b/projects/clr/hipamd/src/hip_memory.cpp index 720d04deb8..aa63bd2c2a 100644 --- a/projects/clr/hipamd/src/hip_memory.cpp +++ b/projects/clr/hipamd/src/hip_memory.cpp @@ -503,24 +503,28 @@ hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocat HIP_INIT_SPECIAL_API(hipArrayCreate, (TRACE_MEM), array, pAllocateArray); HIP_SET_DEVICE(); hipError_t hip_status = hipSuccess; - if (pAllocateArray->width > 0) { + if (pAllocateArray->Width > 0) { auto ctx = ihipGetTlsDefaultCtx(); *array = (hipArray*)malloc(sizeof(hipArray)); - array[0]->drvDesc = *pAllocateArray; - array[0]->width = pAllocateArray->width; - array[0]->height = pAllocateArray->height; + HIP_ARRAY3D_DESCRIPTOR array3D; + array3D.Width = pAllocateArray->Width; + array3D.Height = pAllocateArray->Height; + array3D.Format = pAllocateArray->Format; + array3D.NumChannels = pAllocateArray->NumChannels; + array[0]->width = pAllocateArray->Width; + array[0]->height = pAllocateArray->Height; array[0]->isDrv = true; array[0]->textureType = hipTextureType2D; void** ptr = &array[0]->data; if (ctx) { const unsigned am_flags = 0; - size_t size = pAllocateArray->width; - if (pAllocateArray->height > 0) { - size = size * pAllocateArray->height; + size_t size = pAllocateArray->Width; + if (pAllocateArray->Height > 0) { + size = size * pAllocateArray->Height; } hsa_ext_image_channel_type_t channelType; size_t allocSize = 0; - switch (pAllocateArray->format) { + switch (pAllocateArray->Format) { case HIP_AD_FORMAT_UNSIGNED_INT8: allocSize = size * sizeof(uint8_t); channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8; @@ -568,8 +572,8 @@ hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocat hsa_ext_image_descriptor_t imageDescriptor; - imageDescriptor.width = pAllocateArray->width; - imageDescriptor.height = pAllocateArray->height; + imageDescriptor.width = pAllocateArray->Width; + imageDescriptor.height = pAllocateArray->Height; imageDescriptor.depth = 0; imageDescriptor.array_size = 0; @@ -577,11 +581,11 @@ hipError_t hipArrayCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocat hsa_ext_image_channel_order_t channelOrder; - if (pAllocateArray->numChannels == 4) { + if (pAllocateArray->NumChannels == 4) { channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA; - } else if (pAllocateArray->numChannels == 2) { + } else if (pAllocateArray->NumChannels == 2) { channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG; - } else if (pAllocateArray->numChannels == 1) { + } else if (pAllocateArray->NumChannels == 1) { channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R; } imageDescriptor.format.channel_order = channelOrder; @@ -690,29 +694,28 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, si return ihipLogStatus(hip_status); } -hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAllocateArray) { +hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY3D_DESCRIPTOR* pAllocateArray) { HIP_INIT_SPECIAL_API(hipArray3DCreate, (TRACE_MEM), array, pAllocateArray); hipError_t hip_status = hipSuccess; auto ctx = ihipGetTlsDefaultCtx(); *array = (hipArray*)malloc(sizeof(hipArray)); - array[0]->type = pAllocateArray->flags; - array[0]->width = pAllocateArray->width; - array[0]->height = pAllocateArray->height; - array[0]->depth = pAllocateArray->depth; - array[0]->drvDesc = *pAllocateArray; + array[0]->type = pAllocateArray->Flags; + array[0]->width = pAllocateArray->Width; + array[0]->height = pAllocateArray->Height; + array[0]->depth = pAllocateArray->Depth; array[0]->isDrv = true; array[0]->textureType = hipTextureType3D; void** ptr = &array[0]->data; if (ctx) { const unsigned am_flags = 0; - const size_t size = pAllocateArray->width * pAllocateArray->height * pAllocateArray->depth; + const size_t size = pAllocateArray->Width * pAllocateArray->Height * pAllocateArray->Depth; size_t allocSize = 0; hsa_ext_image_channel_type_t channelType; - switch (pAllocateArray->format) { + switch (pAllocateArray->Format) { case HIP_AD_FORMAT_UNSIGNED_INT8: allocSize = size * sizeof(uint8_t); channelType = HSA_EXT_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8; @@ -760,14 +763,14 @@ hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAlloc &allocGranularity); hsa_ext_image_descriptor_t imageDescriptor; - imageDescriptor.width = pAllocateArray->width; - imageDescriptor.height = pAllocateArray->height; + imageDescriptor.width = pAllocateArray->Width; + imageDescriptor.height = pAllocateArray->Height; imageDescriptor.depth = 0; imageDescriptor.array_size = 0; - switch (pAllocateArray->flags) { + switch (pAllocateArray->Flags) { case hipArrayLayered: imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_2DA; - imageDescriptor.array_size = pAllocateArray->depth; + imageDescriptor.array_size = pAllocateArray->Depth; break; case hipArraySurfaceLoadStore: case hipArrayTextureGather: @@ -777,17 +780,17 @@ hipError_t hipArray3DCreate(hipArray** array, const HIP_ARRAY_DESCRIPTOR* pAlloc case hipArrayCubemap: default: imageDescriptor.geometry = HSA_EXT_IMAGE_GEOMETRY_3D; - imageDescriptor.depth = pAllocateArray->depth; + imageDescriptor.depth = pAllocateArray->Depth; break; } hsa_ext_image_channel_order_t channelOrder; // getChannelOrderAndType(*desc, hipReadModeElementType, &channelOrder, &channelType); - if (pAllocateArray->numChannels == 4) { + if (pAllocateArray->NumChannels == 4) { channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RGBA; - } else if (pAllocateArray->numChannels == 2) { + } else if (pAllocateArray->NumChannels == 2) { channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_RG; - } else if (pAllocateArray->numChannels == 1) { + } else if (pAllocateArray->NumChannels == 1) { channelOrder = HSA_EXT_IMAGE_CHANNEL_ORDER_R; } imageDescriptor.format.channel_order = channelOrder; diff --git a/projects/clr/hipamd/src/hip_texture.cpp b/projects/clr/hipamd/src/hip_texture.cpp index 316fba20cd..e6cd352008 100644 --- a/projects/clr/hipamd/src/hip_texture.cpp +++ b/projects/clr/hipamd/src/hip_texture.cpp @@ -584,7 +584,7 @@ hipError_t ihipBindTextureToArrayImpl(int dim, enum hipTextureReadMode readMode, hsa_ext_image_channel_order_t channelOrder; hsa_ext_image_channel_type_t channelType; if (array->isDrv) { - getDrvChannelOrderAndType(array->drvDesc.format, array->drvDesc.numChannels, + getDrvChannelOrderAndType(array->Format, array->NumChannels, &channelOrder, &channelType); } else { getChannelOrderAndType(desc, readMode, &channelOrder, &channelType); @@ -749,6 +749,6 @@ hipError_t hipTexRefSetAddress2D(textureReference* tex, const HIP_ARRAY_DESCRIPT hipError_t hip_status = hipSuccess; // TODO: hipReadModeElementType is default. hip_status = ihipBindTexture2DImpl(hipTextureType2D, hipReadModeElementType, &offset, devPtr, - NULL, desc->width, desc->height, tex); + NULL, desc->Width, desc->Height, tex); return ihipLogStatus(hip_status); } diff --git a/projects/clr/hipamd/tests/hipify-clang/unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/projects/clr/hipamd/tests/hipify-clang/unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp new file mode 100644 index 0000000000..e7953dc862 --- /dev/null +++ b/projects/clr/hipamd/tests/hipify-clang/unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp @@ -0,0 +1,169 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args +/* +Copyright (c) 2015-present 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. +*/ + +// CHECK: #include +#include +#include +#include +#include + +#define fileName "tex2dKernel.code" +// CHECK: texture tex; +texture tex; +bool testResult = false; + +// CHECK: hipError_t status = cmd; +// CHECK: if (status != hipSuccess) { +// CHECK: std::cout << "error: #" << status << " (" << hipGetErrorString(status) +#define CUDACHECK(cmd) \ + { \ + cudaError_t status = cmd; \ + if (status != cudaSuccess) { \ + std::cout << "error: #" << status << " (" << cudaGetErrorString(status) \ + << ") at line:" << __LINE__ << ": " << #cmd << std::endl; \ + abort(); \ + } \ + } + +bool runTest(int argc, char** argv) { + unsigned int width = 256; + unsigned int height = 256; + unsigned int size = width * height * sizeof(float); + float* hData = (float*)malloc(size); + memset(hData, 0, size); + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + hData[i * width + j] = i * width + j; + } + } + // CHECK: hipModule_t Module; + CUmodule Module; + // CHECK: hipModuleLoad(&Module, fileName); + cuModuleLoad(&Module, fileName); + + // CHECK: hipArray * array; + CUarray array; + // CHECK: HIP_ARRAY_DESCRIPTOR desc; + CUDA_ARRAY_DESCRIPTOR desc; + // CHECK: desc.Format = HIP_AD_FORMAT_FLOAT; + desc.Format = CU_AD_FORMAT_FLOAT; + desc.NumChannels = 1; + desc.Width = width; + desc.Height = height; + // CHECK: hipArrayCreate(&array, &desc); + cuArrayCreate(&array, &desc); + + // CHECK: hip_Memcpy2D copyParam; + CUDA_MEMCPY2D copyParam; + memset(©Param, 0, sizeof(copyParam)); + // CHECK: copyParam.dstMemoryType = hipMemoryTypeArray; + copyParam.dstMemoryType = CU_MEMORYTYPE_ARRAY; + copyParam.dstArray = array; + // CHECK: copyParam.srcMemoryType = hipMemoryTypeHost; + copyParam.srcMemoryType = CU_MEMORYTYPE_HOST; + copyParam.srcHost = hData; + copyParam.srcPitch = width * sizeof(float); + copyParam.WidthInBytes = copyParam.srcPitch; + copyParam.Height = height; + // CHECK: hipMemcpyParam2D(©Param); + cuMemcpy2D(©Param); + + // CHECK: textureReference* texref; + CUtexref_st* texref; + // CHECK: hipModuleGetTexRef(&texref, Module, "tex"); + cuModuleGetTexRef(&texref, Module, "tex"); + // CHECK: hipTexRefSetAddressMode(texref, 0, hipAddressModeWrap); + cuTexRefSetAddressMode(texref, 0, CU_TR_ADDRESS_MODE_WRAP); + // CHECK: hipTexRefSetAddressMode(texref, 1, hipAddressModeWrap); + cuTexRefSetAddressMode(texref, 1, CU_TR_ADDRESS_MODE_WRAP); + // CHECK: hipTexRefSetFilterMode(texref, hipFilterModePoint); + cuTexRefSetFilterMode(texref, CU_TR_FILTER_MODE_POINT); + // CHECK: hipTexRefSetFlags(texref, 0); + cuTexRefSetFlags(texref, 0); + // CHECK: hipTexRefSetFormat(texref, HIP_AD_FORMAT_FLOAT, 1); + cuTexRefSetFormat(texref, CU_AD_FORMAT_FLOAT, 1); + // CHECK: hipTexRefSetArray(texref, array, HIP_TRSA_OVERRIDE_FORMAT); + cuTexRefSetArray(texref, array, CU_TRSA_OVERRIDE_FORMAT); + + float* dData = NULL; + // CHECK: hipMalloc((void**)&dData, size); + cudaMalloc((void**)&dData, size); + + struct { + void* _Ad; + unsigned int _Bd; + unsigned int _Cd; + } args; + args._Ad = (void*) dData; + args._Bd = width; + args._Cd = height; + + size_t sizeTemp = sizeof(args); + + // CHECK: void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, + // CHECK: &sizeTemp, HIP_LAUNCH_PARAM_END}; + void* config[] = {CU_LAUNCH_PARAM_BUFFER_POINTER, &args, CU_LAUNCH_PARAM_BUFFER_SIZE, + &sizeTemp, CU_LAUNCH_PARAM_END}; + + // CHECK: hipFunction_t Function; + CUfunction Function; + // CHECK: hipModuleGetFunction(&Function, Module, "tex2dKernel"); + cuModuleGetFunction(&Function, Module, "tex2dKernel"); + + int temp1 = width / 16; + int temp2 = height / 16; + // CHECK: hipModuleLaunchKernel(Function, 16, 16, 1, temp1, temp2, 1, 0, 0, NULL, (void**)&config); + cuLaunchKernel(Function, 16, 16, 1, temp1, temp2, 1, 0, 0, NULL, (void**)&config); + // CHECK: hipDeviceSynchronize(); + cudaDeviceSynchronize(); + + float* hOutputData = (float*)malloc(size); + memset(hOutputData, 0, size); + // CHECK: hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost); + cudaMemcpy(hOutputData, dData, size, cudaMemcpyDeviceToHost); + + for (int i = 0; i < height; i++) { + for (int j = 0; j < width; j++) { + if (hData[i * width + j] != hOutputData[i * width + j]) { + printf("Difference [ %d %d ]:%f ----%f\n", i, j, hData[i * width + j], + hOutputData[i * width + j]); + testResult = false; + break; + } + } + } + // CHECK: hipFree(dData); + cudaFree(dData); + // CHECK: hipFreeArray(hipArray_t(array)); + cudaFreeArray(cudaArray_t(array)); + return true; +} + +int main(int argc, char** argv) { + // CHECK: hipInit(0); + cuInit(0); + testResult = runTest(argc, argv); + printf("%s ...\n", testResult ? "PASSED" : "FAILED"); + exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE); + return 0; +}