Merge 'master' into 'amd-master'
Change-Id: I473b92d17f663938831469c09f9cc65340ee30f6
This commit is contained in:
Jenkins
+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;
|
||||
|
||||
@@ -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);
|
||||
}
|
||||
Fai riferimento in un nuovo problema
Block a user