AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

[dtest] Adding Test case for hipFuncSetSharedMemConfig() and hipFuncSetAttribute()

Telusuri Sumber 
SWDEV-251504 - QA task for validation

Change-Id: I11ed98f648352cb32161c41290516c77231236bc


[ROCm/hip commit: 548ec20ffd]
This commit is contained in:
Lakhan singh Thakur
2020-10-01 21:33:32 +05:30
committed by Mohan Kumar Mithur
orang tua 3e98b823d3
melakukan 66aba3ea21
2 mengubah file dengan 141 tambahan dan 6 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/tests/src/runtimeApi/module/hipFuncSetSharedMemConfig.cpp
+95 -6
Melihat File
@@ -1,5 +1,5 @@
/*
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2020 - 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
@@ -17,11 +17,100 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip/hip_runtime_api.h>
/* HIT_START
* BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS -std=c++11
* TEST: %t
* HIT_END
*/
// Test Description:
// This test case verifies the working of hipFuncSetSharedMemConfig() api and
// the flag parameter
#include "test_common.h"
int main() {
hipSharedMemConfig_t config;
HIP_PRINT_STATUS(hipFuncSetSharedMemConfig(NULL));
HIP_PRINT_STATUS(hipFuncSetSharedMemConfig(&config));
__global__ void ReverseSeq(int *A, int *B, int N) {
extern __shared__ int SMem[];
int offset = threadIdx.x;
int MirrorVal = N - offset - 1;
SMem[offset] = A[offset];
__syncthreads();
B[offset] = SMem[MirrorVal];
}
int main() {
bool IfTestPassed = true;
int *Ah = NULL, *RAh = NULL, NELMTS = 128;
int *Ad = NULL, *RAd = NULL;
Ah = reinterpret_cast<int*>(malloc(NELMTS * sizeof(int)));
RAh = reinterpret_cast<int*>(malloc(NELMTS * sizeof(int)));
HIPCHECK(hipMalloc(&Ad, NELMTS * sizeof(int)));
HIPCHECK(hipMalloc(&RAd, NELMTS * sizeof(int)));
for (int i = 0; i < NELMTS; ++i) {
Ah[i] = i;
RAh[i] = NELMTS - i - 1;
}
HIPCHECK(hipMemcpy(Ad, Ah, NELMTS * sizeof(int), hipMemcpyHostToDevice));
HIPCHECK(hipMemset(RAd, 0, NELMTS * sizeof(int)));
// Testing hipFuncSetSharedMemConfig() with hipSharedMemBankSizeDefault flag
HIPCHECK(hipFuncSetSharedMemConfig(reinterpret_cast<const void*>(&ReverseSeq),
hipSharedMemBankSizeDefault));
// Kernel Launch with shared mem size of = NELMTS * sizeof(int)
ReverseSeq<<<1, NELMTS, NELMTS * sizeof(int)>>>(Ad, RAd, NELMTS);
memset(Ah, 0, NELMTS * sizeof(int));
// Verifying the results
HIPCHECK(hipMemcpy(Ah, RAd, NELMTS * sizeof(int), hipMemcpyDeviceToHost));
for (int i = 0; i < NELMTS; ++i) {
if (Ah[i] != RAh[i]) {
printf("Mismatch found at %d value of array\n", i);
printf(" after setting the flag hipSharedMemBankSizeDefault\n");
IfTestPassed = false;
}
}
// Testing hipFuncSetSharedMemConfig() with hipSharedMemBankSizeFourBytes flg
HIPCHECK(hipFuncSetSharedMemConfig(reinterpret_cast<const void*>(&ReverseSeq),
hipSharedMemBankSizeFourByte));
HIPCHECK(hipMemset(RAd, 0, NELMTS * sizeof(int)));
// Kernel Launch with shared mem size of = NELMTS * sizeof(int)
ReverseSeq<<<1, NELMTS, NELMTS * sizeof(int)>>>(Ad, RAd, NELMTS);
memset(Ah, 0, NELMTS * sizeof(int));
// Verifying the results
HIPCHECK(hipMemcpy(Ah, RAd, NELMTS * sizeof(int), hipMemcpyDeviceToHost));
for (int i = 0; i < NELMTS; ++i) {
if (Ah[i] != RAh[i]) {
printf("Mismatch found at %d value of array\n", i);
printf(" after setting the flag hipSharedMemBankSizeFourByte\n");
IfTestPassed = false;
}
}
// Testing hipFuncSetSharedMemConfig() with hipSharedMemBankSizeEightBytes flg
HIPCHECK(hipFuncSetSharedMemConfig(reinterpret_cast<const void*>(&ReverseSeq),
hipSharedMemBankSizeEightByte));
HIPCHECK(hipMemset(RAd, 0, NELMTS * sizeof(int)));
// Kernel Launch with shared mem size of = NELMTS * sizeof(int)
ReverseSeq<<<1, NELMTS, NELMTS * sizeof(int)>>>(Ad, RAd, NELMTS);
memset(Ah, 0, NELMTS * sizeof(int));
// Verifying the results
HIPCHECK(hipMemcpy(Ah, RAd, NELMTS * sizeof(int), hipMemcpyDeviceToHost));
for (int i = 0; i < NELMTS; ++i) {
if (Ah[i] != RAh[i]) {
printf("Mismatch found at %d value of array\n", i);
printf(" after setting the flag hipSharedMemBankSizeEightByte\n");
IfTestPassed = false;
}
}
free(Ah);
free(RAh);
HIPCHECK(hipFree(Ad));
HIPCHECK(hipFree(RAd));
if (IfTestPassed) {
passed();
} else {
failed("");
}
}