Added a macro to disable the alignment test as it's broken

[ROCm/hip commit: a76d8f5580]
Этот коммит содержится в:
Srinivasuluch
2018-06-27 19:57:44 +05:30
коммит произвёл GitHub
родитель c2a44fa2e9
Коммит 844928348f
+12 -8
Просмотреть файл
@@ -26,6 +26,8 @@ THE SOFTWARE.
#include "hip/hip_runtime.h"
#include "test_common.h"
#define DISABLE_TEST 0
static const int BLOCK_DIM_SIZE = 1024;
// This test is to verify Struct with variables to check the hipLaunchKernel() support, read and write into the same struct
@@ -69,20 +71,20 @@ typedef struct hipLaunchKernelStruct5 {
typedef struct hipLaunchKernelStruct6 {
char c1;
short int si;
} /*__attribute__ ((aligned(8))) */ hipLaunchKernelStruct_t6;
} __attribute__ ((aligned(8))) hipLaunchKernelStruct_t6;
// This test is to verify struct with aligned(16), right now it's broken on hcc & hip-clang
typedef struct hipLaunchKernelStruct7 {
char c1;
short int si;
} /*__attribute__ ((aligned(16))) */ hipLaunchKernelStruct_t7;
} __attribute__ ((aligned(16))) hipLaunchKernelStruct_t7;
// This test is to verify struct with packed & aligned, size should be 7Bytes, , right now it's broken on hcc & hip-clang
// This test is to verify struct with packed & aligned, size should be 4Bytes, , right now it's broken on hcc & hip-clang
typedef struct hipLaunchKernelStruct8 {
char c1;
short int si;
bool b;
} /* __attribute__ ((packed, aligned(4))) */ hipLaunchKernelStruct_t8;
} __attribute__ ((packed, aligned(4))) hipLaunchKernelStruct_t8;
// Passing struct to a hipLaunchKernel(), read and write into the same struct
__global__ void hipLaunchKernelStructFunc1(hipLaunchParm lp, hipLaunchKernelStruct_t1 hipLaunchKernelStruct_, bool* result_d1) {
@@ -133,7 +135,7 @@ __global__ void hipLaunchKernelStructFunc6(hipLaunchParm lp, hipLaunchKernelStru
// set the result to true if the condition met
int *p = (int*)(&hipLaunchKernelStruct_); // get the address of the struct
// size_t(p)%8 will be 0 if aligned to 8Byte address space
result_d6[x] = ((hipLaunchKernelStruct_.c1 == 'c') && (hipLaunchKernelStruct_.si == 1) /*&& ((size_t(p))%8 ==0)*/) ? true : false;
result_d6[x] = ((hipLaunchKernelStruct_.c1 == 'c') && (hipLaunchKernelStruct_.si == 1) && ((size_t(p))%8 ==0)) ? true : false;
}
// Passing struct which is aligned to 16Byte to a hipLaunchKernel(), set the result_d7 to true if condition met
@@ -143,7 +145,7 @@ __global__ void hipLaunchKernelStructFunc7(hipLaunchParm lp, hipLaunchKernelStru
// set the result to true if the condition met
int *p = (int*)(&hipLaunchKernelStruct_); // get the address of the struct
// size_t(p)%16 will be 0 if aligned to 16Byte address space
result_d7[x] = ((hipLaunchKernelStruct_.c1 == 'c') && (hipLaunchKernelStruct_.si == 1) /*&& ((size_t(p))%16 ==0)*/ ) ? true : false;
result_d7[x] = ((hipLaunchKernelStruct_.c1 == 'c') && (hipLaunchKernelStruct_.si == 1) && ((size_t(p))%16 ==0) ) ? true : false;
}
// Passing struct which is packed & aligned to 4Byte to a hipLaunchKernel(), set the result_d8 to true if condition met
@@ -152,8 +154,8 @@ __global__ void hipLaunchKernelStructFunc8(hipLaunchParm lp, hipLaunchKernelStru
// set the result to true if the condition met
int *p = (int*)(&hipLaunchKernelStruct_); // get the address of the xth element, struct[x],
// size_t(p)%6 will be 0 if aligned to 6Byte address space
result_d8[x] = ((hipLaunchKernelStruct_.c1 == 'c') && (hipLaunchKernelStruct_.si == 1) /*&& ((size_t(p))%4 ==0)*/ ) ? true : false;
// size_t(p)%4 will be 0 if aligned to 4Byte address space
result_d8[x] = ((hipLaunchKernelStruct_.c1 == 'c') && (hipLaunchKernelStruct_.si == 1) && ((size_t(p))%4 ==0) ) ? true : false;
}
__global__ void vAdd(hipLaunchParm lp, float* a) {}
@@ -328,6 +330,7 @@ int main() {
for (int k = 0; k < BLOCK_DIM_SIZE; ++k)
HIPASSERT(result_h5[k] == true);
#if DISABLE_TEST // alignment is broken hence disabled the validation part
// Validation part of the struct, hipLaunchKernelStructFunc6
hipMemcpy(result_h6, result_d6, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost);
for (int k = 0; k < BLOCK_DIM_SIZE; ++k)
@@ -342,6 +345,7 @@ int main() {
hipMemcpy(result_h8, result_d8, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost);
for (int k = 0; k < BLOCK_DIM_SIZE; ++k)
HIPASSERT(result_h8[k] == true);
#endif
// Test case with hipLaunchKernel inside another macro:
float e0;