[HIP][HIPIFY] Split HIP_ARRAY_DESCRIPTOR struct to HIP_ARRAY_DESCRIPTOR and HIP_ARRAY3D_DESCRIPTOR
[Reason] To be compatible with CUDA [#1133]
Update HIP code, hipify-clang, tests and docs
[TODO] Add support of the corresponding functions on nvcc fallback path
[ROCm/clr commit: f0832fd968]
Этот коммит содержится в:
Evgeny Mankov
+11
-11
@@ -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` |
|
||||
|
||||
@@ -27,8 +27,8 @@ const std::map<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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<llvm::StringRef, hipCounter> 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
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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);
|
||||
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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;
|
||||
|
||||
@@ -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);
|
||||
}
|
||||
|
||||
+169
@@ -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 <hip/hip_runtime.h>
|
||||
#include <cuda.h>
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <vector>
|
||||
|
||||
#define fileName "tex2dKernel.code"
|
||||
// CHECK: texture<float, 2, hipReadModeElementType> tex;
|
||||
texture<float, 2, cudaReadModeElementType> 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;
|
||||
}
|
||||
Ссылка в новой задаче
Block a user