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

Update hipLaunchParm.cpp

Browse Source 
wrap up to 80 characters per line
This commit is contained in:
Srinivasuluch
2018-06-27 22:04:12 +05:30
committed by GitHub
parent f69d89e79f
commit 52e808da4d
1 changed files with 101 additions and 77 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/tests/src/kernel/hipLaunchParm.cpp
+101 -77
View File
@@ -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 <sys/time.h>
#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),