Merge remote-tracking branch 'origin/master' into feature_support_globals_for_module_api
This commit is contained in:
Alex Voicu
+11
@@ -8,6 +8,17 @@ We have attempted to document known bugs and limitations - in particular the [HI
|
||||
|
||||
## Revision History:
|
||||
|
||||
===================================================================================================
|
||||
Release: 1.4
|
||||
Date: 2017.10.06
|
||||
- Improvements to HIP event management
|
||||
- Added new HIP_TRACE_API options
|
||||
- Enabled device side assert support
|
||||
- Several bug fixes including hipMallocArray, hipTexture fetch
|
||||
- Support for RHEL/CentOS 7.4
|
||||
- Updates to hipify-perl, hipify-clang and documentation
|
||||
|
||||
|
||||
===================================================================================================
|
||||
Release: 1.3
|
||||
Date: 2017.08.16
|
||||
|
||||
+1
-1
@@ -1,7 +1,7 @@
|
||||
#!/usr/bin/perl -w
|
||||
|
||||
$HIP_BASE_VERSION_MAJOR = "1";
|
||||
$HIP_BASE_VERSION_MINOR = "3";
|
||||
$HIP_BASE_VERSION_MINOR = "4";
|
||||
|
||||
# Need perl > 5.10 to use logic-defined or
|
||||
use 5.006; use v5.10.1;
|
||||
|
||||
@@ -23,6 +23,11 @@ THE SOFTWARE.
|
||||
#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_MATH_FUNCTIONS_H
|
||||
#define HIP_INCLUDE_HIP_HCC_DETAIL_MATH_FUNCTIONS_H
|
||||
|
||||
#if defined(__HCC__)
|
||||
#include <kalmar_math.h>
|
||||
#endif
|
||||
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <hip/hip_vector_types.h>
|
||||
#include <hip/hcc_detail/device_functions.h>
|
||||
|
||||
@@ -51,14 +51,39 @@ hipMemcpyHostToHost
|
||||
,hipMemcpyDefault
|
||||
} hipMemcpyKind ;
|
||||
|
||||
//hipTextureAddressMode
|
||||
#define hipTextureAddressMode cudaTextureAddressMode
|
||||
#define hipAddressModeWrap cudaAddressModeWrap
|
||||
#define hipAddressModeClamp cudaAddressModeClamp
|
||||
#define hipAddressModeMirror cudaAddressModeMirror
|
||||
#define hipAddressModeBorder cudaAddressModeBorder
|
||||
|
||||
//hipTextureFilterMode
|
||||
#define hipTextureFilterMode cudaTextureFilterMode
|
||||
#define hipFilterModePoint cudaFilterModePoint
|
||||
#define hipFilterModeLinear cudaFilterModeLinear
|
||||
|
||||
//hipTextureReadMode
|
||||
#define hipTextureReadMode cudaTextureReadMode
|
||||
#define hipReadModeElementType cudaReadModeElementType
|
||||
#define hipReadModeNormalizedFloat cudaReadModeNormalizedFloat
|
||||
|
||||
typedef enum hipChannelFormatKind {
|
||||
hipChannelFormatKindSigned = 0,
|
||||
hipChannelFormatKindUnsigned = 1,
|
||||
hipChannelFormatKindFloat = 2,
|
||||
hipChannelFormatKindNone = 3
|
||||
}hipChannelFormatKind;
|
||||
|
||||
//hipResourceType
|
||||
#define hipResourceType cudaResourceType
|
||||
#define hipResourceTypeArray cudaResourceTypeArray
|
||||
#define hipResourceTypeMipmappedArray cudaResourceTypeMipmappedArray
|
||||
#define hipResourceTypeLinear cudaResourceTypeLinear
|
||||
#define hipResourceTypePitch2D cudaResourceTypePitch2D
|
||||
//
|
||||
// hipErrorNoDevice.
|
||||
|
||||
/*typedef enum hipTextureFilterMode
|
||||
{
|
||||
hipFilterModePoint = cudaFilterModePoint, ///< Point filter mode.
|
||||
//! @warning cudaFilterModeLinear is not supported.
|
||||
} hipTextureFilterMode;*/
|
||||
#define hipFilterModePoint cudaFilterModePoint
|
||||
|
||||
//! Flags that can be used with hipEventCreateWithFlags:
|
||||
#define hipEventDefault cudaEventDefault
|
||||
@@ -119,18 +144,22 @@ typedef CUdevice hipDevice_t;
|
||||
typedef CUmodule hipModule_t;
|
||||
typedef CUfunction hipFunction_t;
|
||||
typedef CUdeviceptr hipDeviceptr_t;
|
||||
typedef enum cudaChannelFormatKind hipChannelFormatKind;
|
||||
typedef struct cudaChannelFormatDesc hipChannelFormatDesc;
|
||||
typedef enum cudaTextureReadMode hipTextureReadMode;
|
||||
typedef struct cudaArray hipArray;
|
||||
typedef struct cudaArray* hipArray_const_t;
|
||||
#define hipArrayDefault cudaArrayDefault
|
||||
|
||||
typedef cudaTextureObject_t hipTextureObject_t;
|
||||
#define hipTextureType2D cudaTextureType2D;
|
||||
#define hipDeviceMapHost cudaDeviceMapHost
|
||||
|
||||
// Flags that can be used with hipStreamCreateWithFlags
|
||||
#define hipStreamDefault cudaStreamDefault
|
||||
#define hipStreamNonBlocking cudaStreamNonBlocking
|
||||
|
||||
//typedef cudaChannelFormatDesc hipChannelFormatDesc;
|
||||
#define hipChannelFormatDesc cudaChannelFormatDesc
|
||||
|
||||
#define hipResourceDesc cudaResourceDesc
|
||||
#define hipTextureDesc cudaTextureDesc
|
||||
#define hipResourceViewDesc cudaResourceViewDesc
|
||||
//adding code for hipmemSharedConfig
|
||||
#define hipSharedMemBankSizeDefault cudaSharedMemBankSizeDefault
|
||||
#define hipSharedMemBankSizeFourByte cudaSharedMemBankSizeFourByte
|
||||
@@ -214,6 +243,58 @@ inline static enum cudaMemcpyKind hipMemcpyKindToCudaMemcpyKind(hipMemcpyKind ki
|
||||
}
|
||||
}
|
||||
|
||||
inline static cudaTextureAddressMode hipTextureAddressModeToCudaTextureAddressMode(hipTextureAddressMode kind) {
|
||||
switch(kind) {
|
||||
case hipAddressModeWrap:
|
||||
return cudaAddressModeWrap;
|
||||
case hipAddressModeClamp:
|
||||
return cudaAddressModeClamp;
|
||||
case hipAddressModeMirror:
|
||||
return cudaAddressModeMirror;
|
||||
case hipAddressModeBorder:
|
||||
return cudaAddressModeBorder;
|
||||
default:
|
||||
return cudaAddressModeWrap;
|
||||
}
|
||||
}
|
||||
|
||||
inline static cudaTextureFilterMode hipTextureFilterModeToCudaTextureFilterMode(hipTextureFilterMode kind) {
|
||||
switch(kind) {
|
||||
case hipFilterModePoint:
|
||||
return cudaFilterModePoint;
|
||||
case hipFilterModeLinear:
|
||||
return cudaFilterModeLinear;
|
||||
default:
|
||||
return cudaFilterModePoint;
|
||||
}
|
||||
}
|
||||
|
||||
inline static cudaTextureReadMode hipTextureReadModeToCudaTextureReadMode(hipTextureReadMode kind) {
|
||||
switch(kind) {
|
||||
case hipReadModeElementType:
|
||||
return cudaReadModeElementType;
|
||||
case hipReadModeNormalizedFloat:
|
||||
return cudaReadModeNormalizedFloat;
|
||||
default:
|
||||
return cudaReadModeElementType;
|
||||
}
|
||||
}
|
||||
|
||||
inline static cudaChannelFormatKind hipChannelFormatKindToCudaChannelFormatKind(hipChannelFormatKind kind) {
|
||||
switch(kind) {
|
||||
case hipChannelFormatKindSigned :
|
||||
return cudaChannelFormatKindSigned ;
|
||||
case hipChannelFormatKindUnsigned :
|
||||
return cudaChannelFormatKindUnsigned ;
|
||||
case hipChannelFormatKindFloat :
|
||||
return cudaChannelFormatKindFloat ;
|
||||
case hipChannelFormatKindNone :
|
||||
return cudaChannelFormatKindNone ;
|
||||
default:
|
||||
return cudaChannelFormatKindNone ;
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Stream CallBack struct
|
||||
*/
|
||||
@@ -262,8 +343,9 @@ 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 struct hipChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags) {
|
||||
return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags));
|
||||
inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc,
|
||||
size_t width, size_t height, unsigned int flags __dparm(hipArrayDefault)) {
|
||||
return hipCUDAErrorTohipError(cudaMallocArray(array, desc, width, height, flags));
|
||||
}
|
||||
|
||||
inline static hipError_t hipFreeArray(hipArray* array) {
|
||||
@@ -1007,11 +1089,44 @@ inline static hipError_t hipUnbindTexture(struct texture<T, dim, readMode> *tex
|
||||
return hipCUDAErrorTohipError(cudaUnbindTexture(tex));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum hipTextureReadMode readMode>
|
||||
inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode>& tex,
|
||||
hipArray_const_t array,
|
||||
const struct hipChannelFormatDesc& desc)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array,desc));
|
||||
}
|
||||
|
||||
template <class T, int dim, enum cudaTextureReadMode readMode>
|
||||
inline static hipError_t hipBindTextureToArray(struct texture<T, dim, readMode> &tex,
|
||||
hipArray_const_t array)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaBindTextureToArray(tex, array));
|
||||
}
|
||||
|
||||
template <class T>
|
||||
inline static hipChannelFormatDesc hipCreateChannelDesc()
|
||||
{
|
||||
return cudaCreateChannelDesc<T>();
|
||||
}
|
||||
|
||||
inline static hipChannelFormatDesc hipCreateChannelDesc(int x, int y, int z, int w, hipChannelFormatKind f)
|
||||
{
|
||||
return cudaCreateChannelDesc(x , y , z , w, hipChannelFormatKindToCudaChannelFormatKind(f));
|
||||
}
|
||||
|
||||
inline static hipError_t hipCreateTextureObject(hipTextureObject_t* pTexObject,
|
||||
const hipResourceDesc* pResDesc,
|
||||
const hipTextureDesc* pTexDesc,
|
||||
const hipResourceViewDesc* pResViewDesc)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaCreateTextureObject(pTexObject, pResDesc, pTexDesc, pResViewDesc));
|
||||
}
|
||||
|
||||
inline static hipError_t hipDestroyTextureObject(hipTextureObject_t textureObject)
|
||||
{
|
||||
return hipCUDAErrorTohipError(cudaDestroyTextureObject(textureObject));
|
||||
}
|
||||
#endif //__CUDACC__
|
||||
|
||||
#endif //HIP_INCLUDE_HIP_NVCC_DETAIL_HIP_RUNTIME_API_H
|
||||
|
||||
@@ -0,0 +1,112 @@
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp
|
||||
* RUN: %t
|
||||
* HIT_END
|
||||
*/
|
||||
#include <stdlib.h>
|
||||
#include <stdio.h>
|
||||
#include <string.h>
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include "test_common.h"
|
||||
|
||||
bool testResult = true;
|
||||
|
||||
__global__ void tex2DKernel(float* outputData,
|
||||
hipTextureObject_t textureObject,
|
||||
int width,
|
||||
int height)
|
||||
{
|
||||
int x = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x;
|
||||
int y = hipBlockIdx_y*hipBlockDim_y + hipThreadIdx_y;
|
||||
outputData[y*width + x] = tex2D<float>(textureObject, x, y);
|
||||
}
|
||||
|
||||
void runTest(int argc, char **argv);
|
||||
|
||||
int main(int argc, char **argv)
|
||||
{
|
||||
runTest(argc, argv);
|
||||
|
||||
if(testResult) {
|
||||
passed();
|
||||
} else {
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
void 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;
|
||||
}
|
||||
}
|
||||
printf("hData: ");
|
||||
for (int i = 0; i < 64; i++) {
|
||||
printf("%f ", hData[i]);
|
||||
}
|
||||
printf("\n");
|
||||
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat);
|
||||
hipArray *hipArray;
|
||||
hipMallocArray(&hipArray, &channelDesc, width, height);
|
||||
|
||||
hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice);
|
||||
|
||||
struct hipResourceDesc resDesc;
|
||||
memset(&resDesc, 0, sizeof(resDesc));
|
||||
resDesc.resType = hipResourceTypeArray;
|
||||
resDesc.res.array.array = hipArray;
|
||||
|
||||
// Specify texture object parameters
|
||||
struct hipTextureDesc texDesc;
|
||||
memset(&texDesc, 0, sizeof(texDesc));
|
||||
texDesc.addressMode[0] = hipAddressModeWrap;
|
||||
texDesc.addressMode[1] = hipAddressModeWrap;
|
||||
texDesc.filterMode = hipFilterModePoint;
|
||||
texDesc.readMode = hipReadModeElementType;
|
||||
texDesc.normalizedCoords = 0;
|
||||
|
||||
// Create texture object
|
||||
hipTextureObject_t textureObject = 0;
|
||||
hipCreateTextureObject(&textureObject, &resDesc, &texDesc, NULL);
|
||||
|
||||
float* dData = NULL;
|
||||
hipMalloc((void **) &dData, size);
|
||||
|
||||
dim3 dimBlock(16, 16, 1);
|
||||
dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
|
||||
|
||||
hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, textureObject, width, height);
|
||||
|
||||
hipDeviceSynchronize();
|
||||
|
||||
float *hOutputData = (float *) malloc(size);
|
||||
memset(hOutputData, 0, size);
|
||||
hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost);
|
||||
|
||||
printf("dData: ");
|
||||
for (int i = 0; i < 64; i++) {
|
||||
printf("%f ", hOutputData[i]);
|
||||
}
|
||||
printf("\n");
|
||||
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;
|
||||
}
|
||||
}
|
||||
}
|
||||
hipDestroyTextureObject(textureObject);
|
||||
hipFree(dData);
|
||||
hipFreeArray(hipArray);
|
||||
}
|
||||
@@ -0,0 +1,105 @@
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp
|
||||
* RUN: %t
|
||||
* HIT_END
|
||||
*/
|
||||
#include <stdlib.h>
|
||||
#include <stdio.h>
|
||||
#include <string.h>
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include "test_common.h"
|
||||
texture<float, 2, hipReadModeElementType> tex;
|
||||
|
||||
bool testResult = true;
|
||||
|
||||
__global__ void tex2DKernel(float* outputData,
|
||||
hipTextureObject_t textureObject,
|
||||
int width,
|
||||
int height)
|
||||
{
|
||||
int x = hipBlockIdx_x*hipBlockDim_x + hipThreadIdx_x;
|
||||
int y = hipBlockIdx_y*hipBlockDim_y + hipThreadIdx_y;
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
outputData[y*width + x] = tex2D(tex, textureObject, x, y);
|
||||
#else
|
||||
outputData[y*width + x] = tex2D(tex, x, y);
|
||||
#endif
|
||||
}
|
||||
|
||||
void runTest(int argc, char **argv);
|
||||
|
||||
int main(int argc, char **argv)
|
||||
{
|
||||
runTest(argc, argv);
|
||||
if(testResult) {
|
||||
passed();
|
||||
} else {
|
||||
exit(EXIT_FAILURE);
|
||||
}
|
||||
}
|
||||
|
||||
void 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;
|
||||
}
|
||||
}
|
||||
printf("hData: ");
|
||||
for (int i = 0; i < 64; i++) {
|
||||
printf("%f ", hData[i]);
|
||||
}
|
||||
printf("\n");
|
||||
|
||||
hipChannelFormatDesc channelDesc = hipCreateChannelDesc(32, 0, 0, 0, hipChannelFormatKindFloat);
|
||||
hipArray *hipArray;
|
||||
hipMallocArray(&hipArray, &channelDesc, width, height);
|
||||
|
||||
hipMemcpyToArray(hipArray, 0, 0, hData, size, hipMemcpyHostToDevice);
|
||||
|
||||
tex.addressMode[0] = hipAddressModeWrap;
|
||||
tex.addressMode[1] = hipAddressModeWrap;
|
||||
tex.filterMode = hipFilterModePoint;
|
||||
tex.normalized = 0;
|
||||
|
||||
hipBindTextureToArray(tex, hipArray, channelDesc);
|
||||
|
||||
float* dData = NULL;
|
||||
hipMalloc((void **) &dData, size);
|
||||
|
||||
dim3 dimBlock(16, 16, 1);
|
||||
dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1);
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, tex.textureObject, width, height);
|
||||
#else
|
||||
hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, 0, width, height);
|
||||
#endif
|
||||
hipDeviceSynchronize();
|
||||
|
||||
float *hOutputData = (float *) malloc(size);
|
||||
memset(hOutputData, 0, size);
|
||||
hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost);
|
||||
|
||||
printf("dData: ");
|
||||
for (int i = 0; i < 64; i++) {
|
||||
printf("%f ", hOutputData[i]);
|
||||
}
|
||||
printf("\n");
|
||||
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;
|
||||
}
|
||||
}
|
||||
}
|
||||
hipFree(dData);
|
||||
hipFreeArray(hipArray);
|
||||
}
|
||||
Reference in New Issue
Block a user