AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

[dtest] Tests for hipModuleGetTexRef() API

Browse Source 
1. Negative Test Cases
2. API functionality test using multiple streams.
3. API functionality validation using multiple GPUs.
4. Setting hipTexRefSetFlags() with HIP_TRSF_READ_AS_INTEGER flag instead of 0 as suggested in SWDEV-256096.
5. Enabling test cases 0x03 (short) and 0x04 (char).
6. Implemented external code review comments.

SWDEV-238517 for enhancing hip unit tests

Change-Id: If42796047ec1cf2e3695dc2b7f40a2d9dd50f5bd
This commit is contained in:
Rupam Chetia
2020-10-04 20:45:09 +05:30
committed by Mohan Kumar Mithur
parent 44d6914ef3
commit 9d92c90b55
2 changed files with 634 additions and 100 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tests/src/runtimeApi/module/tex2d_kernel.cpp
+40 -11
View File
@@ -19,18 +19,47 @@ 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.
*/
/* HIT_START
* BUILD_CMD: tex2d_kernel.code %hc --genco %S/tex2d_kernel.cpp -o tex2d_kernel.code EXCLUDE_HIP_PLATFORM nvidia EXCLUDE_HIP_RUNTIME rocclr
* HIT_END
*/
#include "hip/hip_runtime.h"
texture<float, 2, hipReadModeElementType> tex;
texture<float, 2, hipReadModeElementType> ftex;
texture<int, 2, hipReadModeElementType> itex;
texture<short, 2, hipReadModeElementType> stex;
texture<char, 2, hipReadModeElementType> ctex;
extern "C" __global__ void tex2dKernel(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);
__device__ float deviceGlobalFloat;
extern "C" __global__ void tex2dKernelFloat(float* outputData,
int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
if ((x < width) && (y < width)) {
outputData[y * width + x] = tex2D(ftex, x, y);
}
}
extern "C" __global__ void tex2dKernelInt(int* outputData,
int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
if ((x < width) && (y < width)) {
outputData[y * width + x] = tex2D(itex, x, y);
}
}
extern "C" __global__ void tex2dKernelInt16(short* outputData,
int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
if ((x < width) && (y < width)) {
outputData[y * width + x] = tex2D(stex, x, y);
}
}
extern "C" __global__ void tex2dKernelInt8(char* outputData,
int width, int height) {
int x = hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x;
int y = hipBlockIdx_y * hipBlockDim_y + hipThreadIdx_y;
if ((x < width) && (y < width)) {
outputData[y * width + x] = tex2D(ctex, x, y);
}
}