diff --git a/hipamd/tests/src/kernel/hipLaunchParm.cpp b/hipamd/tests/src/kernel/hipLaunchParm.cpp index 1568d01dd0..8df10806cb 100644 --- a/hipamd/tests/src/kernel/hipLaunchParm.cpp +++ b/hipamd/tests/src/kernel/hipLaunchParm.cpp @@ -30,7 +30,7 @@ THE SOFTWARE. static const int BLOCK_DIM_SIZE = 1024; -// This test is to verify Struct with variables to check the hipLaunchKernel() +// This test is to verify Struct with variables // support, read from device. typedef struct hipLaunchKernelStruct1 { int li; // local int @@ -58,7 +58,7 @@ typedef struct hipLaunchKernelStruct3 { // This test is to verify empty struct typedef struct hipLaunchKernelStruct4 { - // empty struct, size will be verified from device side, size 1Byte + // empty struct, size will be verified from device side,size 1Byte } hipLaunchKernelStruct_t4; // This test is to verify struct with pointer member variable. @@ -68,22 +68,22 @@ typedef struct hipLaunchKernelStruct5 { } hipLaunchKernelStruct_t5; -// This test is to verify struct with aligned(8), right now it's broken -// on hcc & hip-clang +// This test is to verify struct with aligned(8), +// right now it's broken on hcc & hip-clang typedef struct hipLaunchKernelStruct6 { char c1; short int si; } __attribute__((aligned(8))) hipLaunchKernelStruct_t6; -// This test is to verify struct with aligned(16), right now it's broken -// on hcc & hip-clang +// This test is to verify struct with aligned(16), +// right now it's brokenon hcc & hip-clang typedef struct hipLaunchKernelStruct7 { char c1; short int si; } __attribute__((aligned(16))) hipLaunchKernelStruct_t7; -// This test is to verify struct with packed & aligned, size should be 4Bytes, -// 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; @@ -92,9 +92,10 @@ typedef struct hipLaunchKernelStruct8 { // Passing struct to a hipLaunchKernel(), // read and write into the same struct -__global__ void hipLaunchKernelStructFunc1(hipLaunchParm lp, - hipLaunchKernelStruct_t1 hipLaunchKernelStruct_, - bool* result_d1) { +__global__ void hipLaunchKernelStructFunc1( + hipLaunchParm lp, + hipLaunchKernelStruct_t1 hipLaunchKernelStruct_, + bool* result_d1) { int x = blockIdx.x * blockDim.x + threadIdx.x; // set the result to true if the condition met @@ -105,9 +106,10 @@ __global__ void hipLaunchKernelStructFunc1(hipLaunchParm lp, // Passing struct to a hipLaunchKernel(), checks padding, // read and write into the same struct -__global__ void hipLaunchKernelStructFunc2(hipLaunchParm lp, - hipLaunchKernelStruct_t2 hipLaunchKernelStruct_, - bool* result_d2) { +__global__ void hipLaunchKernelStructFunc2( + hipLaunchParm lp, + hipLaunchKernelStruct_t2 hipLaunchKernelStruct_, + bool* result_d2) { int x = blockIdx.x * blockDim.x + threadIdx.x; // set the result to true if the condition met @@ -119,9 +121,10 @@ __global__ void hipLaunchKernelStructFunc2(hipLaunchParm lp, // Passing struct to a hipLaunchKernel(), checks padding, // read and write into the same struct -__global__ void hipLaunchKernelStructFunc3(hipLaunchParm lp, - hipLaunchKernelStruct_t3 hipLaunchKernelStruct_, - bool* result_d3) { +__global__ void hipLaunchKernelStructFunc3( + hipLaunchParm lp, + hipLaunchKernelStruct_t3 hipLaunchKernelStruct_, + bool* result_d3) { int x = blockIdx.x * blockDim.x + threadIdx.x; // set the result to true if the condition met @@ -131,11 +134,12 @@ __global__ void hipLaunchKernelStructFunc3(hipLaunchParm lp, && (hipLaunchKernelStruct_.bf3 == 1) ); } -// Passing empty struct to a hipLaunchKernel(), check the size of 1Byte, -// set the result_d4 to true if condition met -__global__ void hipLaunchKernelStructFunc4(hipLaunchParm lp, - hipLaunchKernelStruct_t4 hipLaunchKernelStruct_, - bool* result_d4) { +// Passing empty struct to a hipLaunchKernel(), +// check the size of 1Byte, set result_d4 to true if condition met +__global__ void hipLaunchKernelStructFunc4( + hipLaunchParm lp, + hipLaunchKernelStruct_t4 hipLaunchKernelStruct_, + bool* result_d4) { int x = blockIdx.x * blockDim.x + threadIdx.x; // set the result to true if the condition met @@ -143,9 +147,10 @@ __global__ void hipLaunchKernelStructFunc4(hipLaunchParm lp, } // Passing struct with pointer object to a hipLaunchKernel() -__global__ void hipLaunchKernelStructFunc5(hipLaunchParm lp, - hipLaunchKernelStruct_t5 hipLaunchKernelStruct_, - bool* result_d5) { +__global__ void hipLaunchKernelStructFunc5( + hipLaunchParm lp, + hipLaunchKernelStruct_t5 hipLaunchKernelStruct_, + bool* result_d5) { int x = blockIdx.x * blockDim.x + threadIdx.x; // set the result to true if the condition met @@ -155,9 +160,10 @@ __global__ void hipLaunchKernelStructFunc5(hipLaunchParm lp, // Passing struct which is aligned to 8Byte to a hipLaunchKernel(), // set the result_d6 to true if condition met -__global__ void hipLaunchKernelStructFunc6(hipLaunchParm lp, - hipLaunchKernelStruct_t6 hipLaunchKernelStruct_, - bool* result_d6) { +__global__ void hipLaunchKernelStructFunc6( + hipLaunchParm lp, + hipLaunchKernelStruct_t6 hipLaunchKernelStruct_, + bool* result_d6) { int x = blockIdx.x * blockDim.x + threadIdx.x; // set the result to true if the condition met @@ -169,11 +175,12 @@ __global__ void hipLaunchKernelStructFunc6(hipLaunchParm lp, && ((size_t(p))%8 ==0)); } -// Passing struct which is aligned to 16Byte to a hipLaunchKernel(), +// Passing struct which is aligned to 16Byte, // set the result_d7 to true if condition met -__global__ void hipLaunchKernelStructFunc7(hipLaunchParm lp, - hipLaunchKernelStruct_t7 hipLaunchKernelStruct_, - bool* result_d7) { +__global__ void hipLaunchKernelStructFunc7( + hipLaunchParm lp, + hipLaunchKernelStruct_t7 hipLaunchKernelStruct_, + bool* result_d7) { int x = blockIdx.x * blockDim.x + threadIdx.x; // set the result to true if the condition met @@ -185,11 +192,12 @@ __global__ void hipLaunchKernelStructFunc7(hipLaunchParm lp, && ((size_t(p))%16 ==0) ); } -// Passing struct which is packed & aligned to 4Byte to a hipLaunchKernel(), +// Passing struct which is packed & aligned to 4Byte, // set the result_d8 to true if condition met -__global__ void hipLaunchKernelStructFunc8(hipLaunchParm lp, - hipLaunchKernelStruct_t8 hipLaunchKernelStruct_, - bool* result_d8) { +__global__ void hipLaunchKernelStructFunc8( + hipLaunchParm lp, + hipLaunchKernelStruct_t8 hipLaunchKernelStruct_, + bool* result_d8) { int x = blockIdx.x * blockDim.x + threadIdx.x; // set the result to true if the condition met @@ -208,29 +216,29 @@ __global__ void vAdd(hipLaunchParm lp, float* a) {} #define WRAP(...) __VA_ARGS__ #include -#define GPU_PRINT_TIME(cmd, elapsed, quiet) \ - do { \ - struct timeval start, stop; \ - float elapsed; \ - gettimeofday(&start, NULL); \ - hipDeviceSynchronize(); \ - cmd; \ - hipDeviceSynchronize(); \ - gettimeofday(&stop, NULL); \ +#define GPU_PRINT_TIME(cmd, elapsed, quiet) \ + do { \ + struct timeval start, stop; \ + float elapsed; \ + gettimeofday(&start, NULL); \ + hipDeviceSynchronize(); \ + cmd; \ + hipDeviceSynchronize(); \ + gettimeofday(&stop, NULL); \ } while (0); -#define MY_LAUNCH(command, doTrace, msg) \ - { \ - if (doTrace) printf("TRACE: %s %s\n", msg, #command); \ - command; \ +#define MY_LAUNCH(command, doTrace, msg) \ + { \ + if (doTrace) printf("TRACE: %s %s\n", msg, #command); \ + command; \ } -#define MY_LAUNCH_WITH_PAREN(command, doTrace, msg) \ - { \ - if (doTrace) printf("TRACE: %s %s\n", msg, #command); \ - (command); \ +#define MY_LAUNCH_WITH_PAREN(command, doTrace, msg) \ + { \ + if (doTrace) printf("TRACE: %s %s\n", msg, #command); \ + (command); \ } @@ -245,8 +253,9 @@ int main() { hipMalloc((void**)&result_d1, BLOCK_DIM_SIZE*sizeof(bool)); hipHostMalloc((void**)&result_h1, BLOCK_DIM_SIZE*sizeof(bool)); for (int k = 0; k < BLOCK_DIM_SIZE; ++k) { - result_d1[k] = false; // initialize to false, will be set to true - // if the struct is accessible from device. + result_d1[k] = false; + // initialize to false, will be set to + // true if the struct size is 1Byte, from device size } hipLaunchKernelStruct_h1.li = 1; hipLaunchKernelStruct_h1.lf = 1.0; @@ -273,15 +282,17 @@ int main() { hipMalloc((void**)&result_d3, BLOCK_DIM_SIZE*sizeof(bool)); hipHostMalloc((void**)&result_h3, BLOCK_DIM_SIZE*sizeof(bool)); for (int k = 0; k < BLOCK_DIM_SIZE; ++k) { - result_d2[k] = false; // initialize to false, will be set to - // true if the struct is accessible from device. + result_d2[k] = false; + // initialize to false, will be set to + // true if the struct size is 1Byte, from device size } hipLaunchKernelStruct_h3.bf1 = 1; hipLaunchKernelStruct_h3.bf2 = 1; hipLaunchKernelStruct_h3.l1 = 1.0; hipLaunchKernelStruct_h3.bf3 = 1; - hipLaunchKernelStruct_h3.result = false; // This will be set to true if - // the the condition is satisfied, from device side + hipLaunchKernelStruct_h3.result = false; + // initialize to false, will be set to + // true if the struct size is 1Byte, from device size // empty struct hipLaunchKernelStruct_t4 hipLaunchKernelStruct_h4; @@ -289,8 +300,9 @@ int main() { hipMalloc((void**)&result_d4, BLOCK_DIM_SIZE*sizeof(bool)); hipHostMalloc((void**)&result_h4, BLOCK_DIM_SIZE*sizeof(bool)); for (int k = 0; k < BLOCK_DIM_SIZE; ++k) { - result_d4[k] = false; // initialize to false, will be set to - // true if the struct size is 1Byte, from device size + result_d4[k] = false; + // initialize to false, will be set to + // true if the struct size is 1Byte, from device size } // Passing struct with pointer object to a hipLaunchKernel() @@ -306,8 +318,9 @@ int main() { hipLaunchKernelStruct_h5.c1 = 'c'; hipLaunchKernelStruct_h5.cp = cp_d5; for (int k = 0; k < BLOCK_DIM_SIZE; ++k) { - result_d5[k] = false; // initialize to false, will be set to - // true if the struct size is 1Byte, from device size + result_d5[k] = false; + // initialize to false, will be set to + // true if the struct size is 1Byte, from device size } // Passing struct with aligned(8) @@ -318,8 +331,9 @@ int main() { hipLaunchKernelStruct_h6.c1 = 'c'; hipLaunchKernelStruct_h6.si = 1; for (int k = 0; k < BLOCK_DIM_SIZE; ++k) { - result_d6[k] = false; // initialize to false, will be set to - // true if the struct size is 1Byte, from device size + result_d6[k] = false; + // initialize to false, will be set to + // true if the struct size is 1Byte, from device size } // Passing struct with aligned(16) @@ -330,8 +344,9 @@ int main() { hipLaunchKernelStruct_h7.c1 = 'c'; hipLaunchKernelStruct_h7.si = 1; for (int k = 0; k < BLOCK_DIM_SIZE; ++k) { - result_d7[k] = false; // initialize to false, will be set to - // true if the struct size is 1Byte, from device size + result_d7[k] = false; + // initialize to false, will be set to + // true if the struct size is 1Byte, from device size } // Passing struct with packed aligned to 6Bytes hipLaunchKernelStruct_t8 hipLaunchKernelStruct_h8; @@ -341,8 +356,9 @@ int main() { hipLaunchKernelStruct_h8.c1 = 'c'; hipLaunchKernelStruct_h8.si = 1; for (int k = 0; k < BLOCK_DIM_SIZE; ++k) { - result_d8[k] = false; // initialize to false, will be set to - // true if the struct size is 1Byte, from device size + result_d8[k] = false; + // initialize to false, will be set to + // true if the struct size is 1Byte, from device size } // Test the different hipLaunchParm options: @@ -351,21 +367,29 @@ int main() { hipLaunchKernel(vAdd, dim3(1024), 1, 0, 0, Ad); hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad); hipLaunchKernel(hipLaunchKernelStructFunc1, dim3(BLOCK_DIM_SIZE), - dim3(1), 0, 0, hipLaunchKernelStruct_h1, result_d1); + dim3(1), 0, 0, hipLaunchKernelStruct_h1, + result_d1); hipLaunchKernel(hipLaunchKernelStructFunc2, dim3(BLOCK_DIM_SIZE), - dim3(1), 0, 0, hipLaunchKernelStruct_h2, result_d2); + dim3(1), 0, 0, hipLaunchKernelStruct_h2, + result_d2); hipLaunchKernel(hipLaunchKernelStructFunc3, dim3(BLOCK_DIM_SIZE), - dim3(1), 0, 0, hipLaunchKernelStruct_h3, result_d3); + dim3(1), 0, 0, hipLaunchKernelStruct_h3, + result_d3); hipLaunchKernel(hipLaunchKernelStructFunc4, dim3(BLOCK_DIM_SIZE), - dim3(1), 0, 0, hipLaunchKernelStruct_h4, result_d4); + dim3(1), 0, 0, hipLaunchKernelStruct_h4, + result_d4); hipLaunchKernel(hipLaunchKernelStructFunc5, dim3(BLOCK_DIM_SIZE), - dim3(1), 0, 0, hipLaunchKernelStruct_h5, result_d5); + dim3(1), 0, 0, hipLaunchKernelStruct_h5, + result_d5); hipLaunchKernel(hipLaunchKernelStructFunc6, dim3(BLOCK_DIM_SIZE), - dim3(1), 0, 0, hipLaunchKernelStruct_h6, result_d6); + dim3(1), 0, 0, hipLaunchKernelStruct_h6, + result_d6); hipLaunchKernel(hipLaunchKernelStructFunc7, dim3(BLOCK_DIM_SIZE), - dim3(1), 0, 0, hipLaunchKernelStruct_h7, result_d7); + dim3(1), 0, 0, hipLaunchKernelStruct_h7, + result_d7); hipLaunchKernel(hipLaunchKernelStructFunc8, dim3(BLOCK_DIM_SIZE), - dim3(1), 0, 0, hipLaunchKernelStruct_h8, result_d8); + dim3(1), 0, 0, hipLaunchKernelStruct_h8, + result_d8); // Validation part of the struct, hipLaunchKernelStructFunc1 hipMemcpy(result_h1, result_d1, BLOCK_DIM_SIZE*sizeof(bool),