Added texture 2D driver API usage example
Этот коммит содержится в:
Rahul Garg
@@ -0,0 +1,17 @@
|
||||
HIP_PATH?= $(wildcard /opt/rocm/hip)
|
||||
ifeq (,$(HIP_PATH))
|
||||
HIP_PATH=../../..
|
||||
endif
|
||||
HIPCC=$(HIP_PATH)/bin/hipcc
|
||||
HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler)
|
||||
|
||||
all: tex2dKernel.code texture2dDrv.out
|
||||
|
||||
texture2dDrv.out: texture2dDrv.cpp
|
||||
$(HIPCC) $(HIPCC_FLAGS) $< -o $@
|
||||
|
||||
tex2dKernel.code: tex2dKernel.cpp
|
||||
$(HIPCC) --genco $(GENCO_FLAGS) $^ -o $@
|
||||
|
||||
clean:
|
||||
rm -f *.code *.out
|
||||
@@ -0,0 +1,33 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
extern texture<float, 2, hipReadModeElementType> tex;
|
||||
|
||||
__global__ void tex2dKernel(hipLaunchParm lp, float* outputData,
|
||||
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(tex, x, y);
|
||||
}
|
||||
@@ -0,0 +1,153 @@
|
||||
/*
|
||||
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.
|
||||
*/
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip/hip_runtime_api.h"
|
||||
#include <iostream>
|
||||
#include <fstream>
|
||||
#include <vector>
|
||||
#include <hip/hip_hcc.h>
|
||||
|
||||
#define fileName "tex2dKernel.code"
|
||||
|
||||
texture<float, 2, hipReadModeElementType> tex;
|
||||
bool testResult = true;
|
||||
|
||||
#define HIP_CHECK(cmd) \
|
||||
{\
|
||||
hipError_t status = cmd;\
|
||||
if(status != hipSuccess) {std::cout<<"error: #"<<status<<" ("<< hipGetErrorString(status) << ") at line:"<<__LINE__<<": "<<#cmd<<std::endl;abort();}\
|
||||
}
|
||||
|
||||
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;
|
||||
}
|
||||
}
|
||||
hipModule_t Module;
|
||||
HIP_CHECK(hipModuleLoad(&Module, fileName));
|
||||
|
||||
hipArray* array;
|
||||
HIP_ARRAY_DESCRIPTOR desc;
|
||||
desc.format = HIP_AD_FORMAT_FLOAT;
|
||||
desc.numChannels = 1;
|
||||
desc.width = width;
|
||||
desc.height = height;
|
||||
hipArrayCreate(&array, &desc);
|
||||
|
||||
hip_Memcpy2D copyParam;
|
||||
memset(©Param, 0, sizeof(copyParam));
|
||||
copyParam.dstMemoryType = hipMemoryTypeArray;
|
||||
copyParam.dstArray = array;
|
||||
copyParam.srcMemoryType = hipMemoryTypeHost;
|
||||
copyParam.srcHost = hData;
|
||||
copyParam.srcPitch = width * sizeof(float);
|
||||
copyParam.widthInBytes = copyParam.srcPitch;
|
||||
copyParam.height = height;
|
||||
hipMemcpy_2D(©Param);
|
||||
|
||||
hipTexRefSetAddressMode(&tex, 0, hipAddressModeWrap);
|
||||
hipTexRefSetAddressMode(&tex, 1, hipAddressModeWrap);
|
||||
hipTexRefSetFilterMode(&tex, hipFilterModePoint);
|
||||
hipTexRefSetFlags(&tex, 0);
|
||||
hipTexRefSetFormat(&tex, HIP_AD_FORMAT_FLOAT, 1);
|
||||
hipTexRefSetArray(&tex, array, HIP_TRSA_OVERRIDE_FORMAT);
|
||||
|
||||
float* dData = NULL;
|
||||
hipMalloc((void **) &dData, size);
|
||||
|
||||
#ifdef __HIP_PLATFORM_HCC__
|
||||
|
||||
struct {
|
||||
uint32_t _hidden[6]; // genco path + wrapper-gen pass used hidden arguments.
|
||||
void * _Ad;
|
||||
unsigned int _Bd;
|
||||
unsigned int _Cd;
|
||||
} args;
|
||||
args._Ad = dData;
|
||||
args._Bd = width;
|
||||
args._Cd = height;
|
||||
|
||||
#endif
|
||||
|
||||
#ifdef __HIP_PLATFORM_NVCC__
|
||||
struct {
|
||||
uint32_t _hidden[1];
|
||||
void * _Ad;
|
||||
unsigned int _Bd;
|
||||
unsigned int _Cd;
|
||||
} args;
|
||||
|
||||
args._hidden[0] = 0;
|
||||
args._Ad = dData;
|
||||
args._Bd = width;
|
||||
args._Cd = height;
|
||||
#endif
|
||||
|
||||
|
||||
size_t sizeTemp = sizeof(args);
|
||||
|
||||
void *config[] = {
|
||||
HIP_LAUNCH_PARAM_BUFFER_POINTER, &args,
|
||||
HIP_LAUNCH_PARAM_BUFFER_SIZE, &sizeTemp,
|
||||
HIP_LAUNCH_PARAM_END
|
||||
};
|
||||
|
||||
hipFunction_t Function;
|
||||
HIP_CHECK(hipModuleGetFunction(&Function, Module, "tex2dKernel"));
|
||||
|
||||
int temp1= width/16;
|
||||
int temp2 = height/16;
|
||||
HIP_CHECK(hipModuleLaunchKernel(Function, 16, 16, 1, temp1, temp2, 1, 0, 0, NULL, (void**)&config));
|
||||
hipDeviceSynchronize();
|
||||
|
||||
float *hOutputData = (float *) malloc(size);
|
||||
memset(hOutputData, 0, size);
|
||||
hipMemcpy(hOutputData, dData, size, hipMemcpyDeviceToHost);
|
||||
|
||||
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(array);
|
||||
}
|
||||
|
||||
int main(int argc, char **argv){
|
||||
hipInit(0);
|
||||
runTest(argc, argv);
|
||||
printf("%s ...\n", testResult ? "PASSED" : "FAILED");
|
||||
exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE);
|
||||
return 0;
|
||||
}
|
||||
Ссылка в новой задаче
Block a user