From 0e98aa4c7869ce1af0d5a773fc1edc90798d7c8b Mon Sep 17 00:00:00 2001 From: srinivas Charupally Date: Mon, 11 Jun 2018 12:37:38 +0530 Subject: [PATCH 01/12] Adding struct test for hipLaunchKernel() --- tests/src/kernel/hipLaunchParm.cpp | 11 ++++++++++- 1 file changed, 10 insertions(+), 1 deletion(-) diff --git a/tests/src/kernel/hipLaunchParm.cpp b/tests/src/kernel/hipLaunchParm.cpp index 46734a7b6b..45d50aa713 100644 --- a/tests/src/kernel/hipLaunchParm.cpp +++ b/tests/src/kernel/hipLaunchParm.cpp @@ -28,8 +28,14 @@ THE SOFTWARE. #include "hip/hip_runtime_api.h" #include -__global__ void vAdd(hipLaunchParm lp, float* a) {} +// Struct with dummy variables to check the hipLaunchKernel() +typedef struct hipLaunchKernelStruct { + int i; + float f; +} hipLaunchKernelStruct_t; +__global__ void vAdd(hipLaunchParm lp, float* a) {} +__global__ void hipLaunchKernelStructFunc(hipLaunchParm lp, hipLaunchKernelStruct_t* hipLaunchKernelStruct_) {} //--- // Some wrapper macro for testing: @@ -64,13 +70,16 @@ __global__ void vAdd(hipLaunchParm lp, float* a) {} int main() { float* Ad; + hipLaunchKernelStruct_t* hipLaunchKernelStruct_; hipMalloc((void**)&Ad, 1024); + hipMalloc((void**)&hipLaunchKernelStruct_, 1024); // Test the different hipLaunchParm options: hipLaunchKernel(vAdd, size_t(1024), 1, 0, 0, Ad); hipLaunchKernel(vAdd, 1024, dim3(1), 0, 0, Ad); hipLaunchKernel(vAdd, dim3(1024), 1, 0, 0, Ad); hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad); + hipLaunchKernel(hipLaunchKernelStructFunc, dim3(1024), dim3(1), 0, 0, hipLaunchKernelStruct_); // Test case with hipLaunchKernel inside another macro: float e0; From ab07d7dbe6d97df806f8ed66101011a6c2b42731 Mon Sep 17 00:00:00 2001 From: Srinivasuluch Date: Mon, 11 Jun 2018 15:21:20 +0530 Subject: [PATCH 02/12] Update hipLaunchParm.cpp cleanup --- tests/src/kernel/hipLaunchParm.cpp | 2 -- 1 file changed, 2 deletions(-) diff --git a/tests/src/kernel/hipLaunchParm.cpp b/tests/src/kernel/hipLaunchParm.cpp index 45d50aa713..ce3004dcd4 100644 --- a/tests/src/kernel/hipLaunchParm.cpp +++ b/tests/src/kernel/hipLaunchParm.cpp @@ -25,8 +25,6 @@ THE SOFTWARE. #include "hip/hip_runtime.h" #include "test_common.h" -#include "hip/hip_runtime_api.h" -#include // Struct with dummy variables to check the hipLaunchKernel() typedef struct hipLaunchKernelStruct { From d8290f5a7bf6cb26233a463511d87627e8f1bc64 Mon Sep 17 00:00:00 2001 From: Srinivasuluch Date: Mon, 25 Jun 2018 16:53:12 +0530 Subject: [PATCH 03/12] Update hipLaunchParm.cpp Added validation part for struct, added two struct padding tests --- tests/src/kernel/hipLaunchParm.cpp | 114 ++++++++++++++++++++++++++--- 1 file changed, 105 insertions(+), 9 deletions(-) diff --git a/tests/src/kernel/hipLaunchParm.cpp b/tests/src/kernel/hipLaunchParm.cpp index ce3004dcd4..b0f03115c5 100644 --- a/tests/src/kernel/hipLaunchParm.cpp +++ b/tests/src/kernel/hipLaunchParm.cpp @@ -26,14 +26,60 @@ THE SOFTWARE. #include "hip/hip_runtime.h" #include "test_common.h" -// Struct with dummy variables to check the hipLaunchKernel() -typedef struct hipLaunchKernelStruct { - int i; - float f; -} hipLaunchKernelStruct_t; +static const int STRUCT_SIZE = 1024; + +// This test is to verify Struct with variables to check the hipLaunchKernel() support, read and write into the same struct +typedef struct hipLaunchKernelStruct1 { + int li; // local int + float lf; // local float + bool result; // default is false, will be set to true if the condition is met +} hipLaunchKernelStruct_t1; + +// This test is to verify struct with padding, read and write into the same struct +typedef struct hipLaunchKernelStruct2 { + char c1; // local char + long l1; // local long + char c2; // local char + long l2; // local long + bool result; // default is false, will be set to true if the condition is met +} hipLaunchKernelStruct_t2; + +typedef struct hipLaunchKernelStruct3 { + char bf1; + char bf2; + long l1; + char bf3; + bool result; // default is false, will be set to true if the condition is met +} hipLaunchKernelStruct_t3; + + +// Passing struct to a hipLaunchKernel(), read and write into the same struct +__global__ void hipLaunchKernelStructFunc1(hipLaunchParm lp, hipLaunchKernelStruct_t1* hipLaunchKernelStruct_) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + hipLaunchKernelStruct_[x].result = ((hipLaunchKernelStruct_[x].li == 1) && (hipLaunchKernelStruct_[x].lf == 1.0)) ? true : false; +} + +// Passing struct to a hipLaunchKernel(), checks padding, read and write into the same struct +__global__ void hipLaunchKernelStructFunc2(hipLaunchParm lp, hipLaunchKernelStruct_t2* hipLaunchKernelStruct_) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + hipLaunchKernelStruct_[x].result = ((hipLaunchKernelStruct_[x].c1 == 'a') && (hipLaunchKernelStruct_[x].l1 == 1.0) + && (hipLaunchKernelStruct_[x].c2 == 'b') && (hipLaunchKernelStruct_[x].l2 == 2.0) ) ? true : false; +} + +// Passing struct to a hipLaunchKernel(), checks padding, read and write into the same struct +__global__ void hipLaunchKernelStructFunc3(hipLaunchParm lp, hipLaunchKernelStruct_t3* hipLaunchKernelStruct_) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + hipLaunchKernelStruct_[x].result = ((hipLaunchKernelStruct_[x].bf1 == 1) && (hipLaunchKernelStruct_[x].bf2 == 1) + && (hipLaunchKernelStruct_[x].l1 == 1.0) && (hipLaunchKernelStruct_[x].bf3 == 1) ) ? true : false; +} __global__ void vAdd(hipLaunchParm lp, float* a) {} -__global__ void hipLaunchKernelStructFunc(hipLaunchParm lp, hipLaunchKernelStruct_t* hipLaunchKernelStruct_) {} //--- // Some wrapper macro for testing: @@ -68,16 +114,66 @@ __global__ void hipLaunchKernelStructFunc(hipLaunchParm lp, hipLaunchKernelStruc int main() { float* Ad; - hipLaunchKernelStruct_t* hipLaunchKernelStruct_; + hipMalloc((void**)&Ad, 1024); - hipMalloc((void**)&hipLaunchKernelStruct_, 1024); + + // Struct type, check access from device. + hipLaunchKernelStruct_t1 *hipLaunchKernelStruct_d1, *hipLaunchKernelStruct_h1; + hipMalloc((void**)&hipLaunchKernelStruct_d1, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t1)); + hipHostMalloc((void**)&hipLaunchKernelStruct_h1, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t1)); + for (int k = 0; k < STRUCT_SIZE; ++k) { + hipLaunchKernelStruct_d1[k].li = 1; + hipLaunchKernelStruct_d1[k].lf = 1.0; + hipLaunchKernelStruct_d1[k].result = false; // This will be set to true if the the condition is satisfied, from device side + } + + // Struct type, checks padding + hipLaunchKernelStruct_t2 *hipLaunchKernelStruct_d2, *hipLaunchKernelStruct_h2; + hipMalloc((void**)&hipLaunchKernelStruct_d2, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t2)); + hipHostMalloc((void**)&hipLaunchKernelStruct_h2, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t2)); + for (int k = 0; k < STRUCT_SIZE; ++k) { + hipLaunchKernelStruct_d2[k].c1 = 'a'; + hipLaunchKernelStruct_d2[k].l1 = 1.0; + hipLaunchKernelStruct_d2[k].c2 = 'b'; + hipLaunchKernelStruct_d2[k].l2 = 2.0; + hipLaunchKernelStruct_d2[k].result = false; // This will be set to true if the the condition is satisfied, from device side + } + + // Struct type, checks padding, assigning integer to a char + hipLaunchKernelStruct_t3 *hipLaunchKernelStruct_d3, *hipLaunchKernelStruct_h3; + hipMalloc((void**)&hipLaunchKernelStruct_d3, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t3)); + hipHostMalloc((void**)&hipLaunchKernelStruct_h3, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t3)); + for (int k = 0; k < STRUCT_SIZE; ++k) { + hipLaunchKernelStruct_d3[k].bf1 = 1; + hipLaunchKernelStruct_d3[k].bf2 = 1; + hipLaunchKernelStruct_d3[k].l1 = 1.0; + hipLaunchKernelStruct_d3[k].bf3 = 1; + hipLaunchKernelStruct_d3[k].result = false; // This will be set to true if the the condition is satisfied, from device side + } // Test the different hipLaunchParm options: hipLaunchKernel(vAdd, size_t(1024), 1, 0, 0, Ad); hipLaunchKernel(vAdd, 1024, dim3(1), 0, 0, Ad); hipLaunchKernel(vAdd, dim3(1024), 1, 0, 0, Ad); hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad); - hipLaunchKernel(hipLaunchKernelStructFunc, dim3(1024), dim3(1), 0, 0, hipLaunchKernelStruct_); + hipLaunchKernel(hipLaunchKernelStructFunc1, dim3(STRUCT_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_d1); + hipLaunchKernel(hipLaunchKernelStructFunc2, dim3(STRUCT_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_d2); + hipLaunchKernel(hipLaunchKernelStructFunc3, dim3(STRUCT_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_d3); + + // Validation part of the struct + hipMemcpy(hipLaunchKernelStruct_h1, hipLaunchKernelStruct_d1, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t1), hipMemcpyDeviceToHost); + for (int k = 0; k < STRUCT_SIZE; ++k) + HIPASSERT(hipLaunchKernelStruct_h1[k].result == true); + + // Validation part of the struct + hipMemcpy(hipLaunchKernelStruct_h2, hipLaunchKernelStruct_d2, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t2), hipMemcpyDeviceToHost); + for (int k = 0; k < STRUCT_SIZE; ++k) + HIPASSERT(hipLaunchKernelStruct_h2[k].result == true); + + // Validation part of the struct + hipMemcpy(hipLaunchKernelStruct_h3, hipLaunchKernelStruct_d3, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t3), hipMemcpyDeviceToHost); + for (int k = 0; k < STRUCT_SIZE; ++k) + HIPASSERT(hipLaunchKernelStruct_h3[k].result == true); // Test case with hipLaunchKernel inside another macro: float e0; From 6bd047fc47402c4a65aea2d31a0cc0fff84d86bf Mon Sep 17 00:00:00 2001 From: Srinivasuluch Date: Wed, 27 Jun 2018 16:45:51 +0530 Subject: [PATCH 04/12] Update hipLaunchParm.cpp pass by value for struct --- tests/src/kernel/hipLaunchParm.cpp | 296 +++++++++++++++++++++++------ 1 file changed, 240 insertions(+), 56 deletions(-) diff --git a/tests/src/kernel/hipLaunchParm.cpp b/tests/src/kernel/hipLaunchParm.cpp index b0f03115c5..2b8a29c329 100644 --- a/tests/src/kernel/hipLaunchParm.cpp +++ b/tests/src/kernel/hipLaunchParm.cpp @@ -26,57 +26,134 @@ THE SOFTWARE. #include "hip/hip_runtime.h" #include "test_common.h" -static const int STRUCT_SIZE = 1024; +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 typedef struct hipLaunchKernelStruct1 { int li; // local int float lf; // local float - bool result; // default is false, will be set to true if the condition is met + bool result; // local bool } hipLaunchKernelStruct_t1; // This test is to verify struct with padding, read and write into the same struct typedef struct hipLaunchKernelStruct2 { - char c1; // local char - long l1; // local long - char c2; // local char - long l2; // local long - bool result; // default is false, will be set to true if the condition is met + char c1; + long l1; + char c2; + long l2; + bool result; } hipLaunchKernelStruct_t2; +// This test is to verify struct with padding, read and write into the same struct typedef struct hipLaunchKernelStruct3 { char bf1; char bf2; long l1; char bf3; - bool result; // default is false, will be set to true if the condition is met + bool result; } hipLaunchKernelStruct_t3; +// This test is to verify empty struct +typedef struct hipLaunchKernelStruct4 { + // empty struct, size will be verified from device side, size 1Byte +} hipLaunchKernelStruct_t4; + +// This test is to verify struct with pointer member variable. +typedef struct hipLaunchKernelStruct5 { + char c1; + char* cp; // char pointer +} hipLaunchKernelStruct_t5; + + +// 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 +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 7Bytes, , 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; // Passing struct to a hipLaunchKernel(), read and write into the same struct -__global__ void hipLaunchKernelStructFunc1(hipLaunchParm lp, hipLaunchKernelStruct_t1* hipLaunchKernelStruct_) { +__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 - hipLaunchKernelStruct_[x].result = ((hipLaunchKernelStruct_[x].li == 1) && (hipLaunchKernelStruct_[x].lf == 1.0)) ? true : false; + result_d1[x] = ((hipLaunchKernelStruct_.li == 1) && (hipLaunchKernelStruct_.lf == 1.0) && (hipLaunchKernelStruct_.result == false)) ? true : false; } // Passing struct to a hipLaunchKernel(), checks padding, read and write into the same struct -__global__ void hipLaunchKernelStructFunc2(hipLaunchParm lp, hipLaunchKernelStruct_t2* hipLaunchKernelStruct_) { +__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 - hipLaunchKernelStruct_[x].result = ((hipLaunchKernelStruct_[x].c1 == 'a') && (hipLaunchKernelStruct_[x].l1 == 1.0) - && (hipLaunchKernelStruct_[x].c2 == 'b') && (hipLaunchKernelStruct_[x].l2 == 2.0) ) ? true : false; + result_d2[x] = ((hipLaunchKernelStruct_.c1 == 'a') && (hipLaunchKernelStruct_.l1 == 1.0) + && (hipLaunchKernelStruct_.c2 == 'b') && (hipLaunchKernelStruct_.l2 == 2.0) ) ? true : false; } // Passing struct to a hipLaunchKernel(), checks padding, read and write into the same struct -__global__ void hipLaunchKernelStructFunc3(hipLaunchParm lp, hipLaunchKernelStruct_t3* hipLaunchKernelStruct_) { +__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 - hipLaunchKernelStruct_[x].result = ((hipLaunchKernelStruct_[x].bf1 == 1) && (hipLaunchKernelStruct_[x].bf2 == 1) - && (hipLaunchKernelStruct_[x].l1 == 1.0) && (hipLaunchKernelStruct_[x].bf3 == 1) ) ? true : false; + result_d3[x] = ((hipLaunchKernelStruct_.bf1 == 1) && (hipLaunchKernelStruct_.bf2 == 1) + && (hipLaunchKernelStruct_.l1 == 1.0) && (hipLaunchKernelStruct_.bf3 == 1) ) ? true : false; +} + +// 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) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + result_d4[x] = (sizeof(hipLaunchKernelStruct_) == 1) ? true : false; +} + +// Passing struct with pointer object to a hipLaunchKernel() +__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 + result_d5[x] = ((hipLaunchKernelStruct_.c1 == 'c') && (*hipLaunchKernelStruct_.cp == 'p')) ? true : false; +} + +// 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) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // 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; +} + +// Passing struct which is aligned to 16Byte to a hipLaunchKernel(), set the result_d7 to true if condition met +__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 + 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; +} + +// Passing struct which is packed & aligned to 4Byte to a hipLaunchKernel(), set the result_d8 to true if condition met +__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 + 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; } __global__ void vAdd(hipLaunchParm lp, float* a) {} @@ -118,37 +195,98 @@ int main() { hipMalloc((void**)&Ad, 1024); // Struct type, check access from device. - hipLaunchKernelStruct_t1 *hipLaunchKernelStruct_d1, *hipLaunchKernelStruct_h1; - hipMalloc((void**)&hipLaunchKernelStruct_d1, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t1)); - hipHostMalloc((void**)&hipLaunchKernelStruct_h1, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t1)); - for (int k = 0; k < STRUCT_SIZE; ++k) { - hipLaunchKernelStruct_d1[k].li = 1; - hipLaunchKernelStruct_d1[k].lf = 1.0; - hipLaunchKernelStruct_d1[k].result = false; // This will be set to true if the the condition is satisfied, from device side + hipLaunchKernelStruct_t1 hipLaunchKernelStruct_h1; + bool *result_d1, *result_h1; + 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. } + hipLaunchKernelStruct_h1.li = 1; + hipLaunchKernelStruct_h1.lf = 1.0; + hipLaunchKernelStruct_h1.result = false; // Struct type, checks padding - hipLaunchKernelStruct_t2 *hipLaunchKernelStruct_d2, *hipLaunchKernelStruct_h2; - hipMalloc((void**)&hipLaunchKernelStruct_d2, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t2)); - hipHostMalloc((void**)&hipLaunchKernelStruct_h2, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t2)); - for (int k = 0; k < STRUCT_SIZE; ++k) { - hipLaunchKernelStruct_d2[k].c1 = 'a'; - hipLaunchKernelStruct_d2[k].l1 = 1.0; - hipLaunchKernelStruct_d2[k].c2 = 'b'; - hipLaunchKernelStruct_d2[k].l2 = 2.0; - hipLaunchKernelStruct_d2[k].result = false; // This will be set to true if the the condition is satisfied, from device side + hipLaunchKernelStruct_t2 hipLaunchKernelStruct_h2; + bool *result_d2, *result_h2; + hipMalloc((void**)&result_d2, BLOCK_DIM_SIZE*sizeof(bool)); + hipHostMalloc((void**)&result_h2, 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. } + hipLaunchKernelStruct_h2.c1 = 'a'; + hipLaunchKernelStruct_h2.l1 = 1.0; + hipLaunchKernelStruct_h2.c2 = 'b'; + hipLaunchKernelStruct_h2.l2 = 2.0; + hipLaunchKernelStruct_h2.result = false; // Struct type, checks padding, assigning integer to a char - hipLaunchKernelStruct_t3 *hipLaunchKernelStruct_d3, *hipLaunchKernelStruct_h3; - hipMalloc((void**)&hipLaunchKernelStruct_d3, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t3)); - hipHostMalloc((void**)&hipLaunchKernelStruct_h3, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t3)); - for (int k = 0; k < STRUCT_SIZE; ++k) { - hipLaunchKernelStruct_d3[k].bf1 = 1; - hipLaunchKernelStruct_d3[k].bf2 = 1; - hipLaunchKernelStruct_d3[k].l1 = 1.0; - hipLaunchKernelStruct_d3[k].bf3 = 1; - hipLaunchKernelStruct_d3[k].result = false; // This will be set to true if the the condition is satisfied, from device side + hipLaunchKernelStruct_t3 hipLaunchKernelStruct_h3; + bool *result_d3, *result_h3; + 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. + } + 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 + + // empty struct + hipLaunchKernelStruct_t4 hipLaunchKernelStruct_h4; + bool *result_d4, *result_h4; + 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 + } + + // Passing struct with pointer object to a hipLaunchKernel() + hipLaunchKernelStruct_t5 hipLaunchKernelStruct_h5; + char* cp_d5; // This is passed as pointer to struct member, struct.cp = &cp_d5 + bool *result_d5, *result_h5; + hipMalloc((void**)&result_d5, BLOCK_DIM_SIZE*sizeof(bool)); + hipMalloc((void**)&cp_d5, sizeof(char)); // allocating memory for char pointer on device + hipHostMalloc((void**)&result_h5, BLOCK_DIM_SIZE*sizeof(bool)); + *cp_d5 = 'p'; // initializing memory to 'p' + 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 + } + + // Passing struct with aligned(8) + hipLaunchKernelStruct_t6 hipLaunchKernelStruct_h6; + bool *result_d6, *result_h6; + hipMalloc((void**)&result_d6, BLOCK_DIM_SIZE*sizeof(bool)); + hipHostMalloc((void**)&result_h6, BLOCK_DIM_SIZE*sizeof(bool)); + 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 + } + + // Passing struct with aligned(16) + hipLaunchKernelStruct_t7 hipLaunchKernelStruct_h7; + bool *result_d7, *result_h7; + hipMalloc((void**)&result_d7, BLOCK_DIM_SIZE*sizeof(bool)); + hipHostMalloc((void**)&result_h7, BLOCK_DIM_SIZE*sizeof(bool)); + 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 + } + // Passing struct with packed aligned to 6Bytes + hipLaunchKernelStruct_t8 hipLaunchKernelStruct_h8; + bool *result_d8, *result_h8; + hipMalloc((void**)&result_d8, BLOCK_DIM_SIZE*sizeof(bool)); + hipHostMalloc((void**)&result_h8, BLOCK_DIM_SIZE*sizeof(bool)); + 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 } // Test the different hipLaunchParm options: @@ -156,24 +294,54 @@ int main() { hipLaunchKernel(vAdd, 1024, dim3(1), 0, 0, Ad); hipLaunchKernel(vAdd, dim3(1024), 1, 0, 0, Ad); hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad); - hipLaunchKernel(hipLaunchKernelStructFunc1, dim3(STRUCT_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_d1); - hipLaunchKernel(hipLaunchKernelStructFunc2, dim3(STRUCT_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_d2); - hipLaunchKernel(hipLaunchKernelStructFunc3, dim3(STRUCT_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_d3); + hipLaunchKernel(hipLaunchKernelStructFunc1, dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h1, result_d1); + hipLaunchKernel(hipLaunchKernelStructFunc2, dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h2, result_d2); + hipLaunchKernel(hipLaunchKernelStructFunc3, dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h3, result_d3); + hipLaunchKernel(hipLaunchKernelStructFunc4, dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h4, result_d4); + hipLaunchKernel(hipLaunchKernelStructFunc5, dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h5, result_d5); + hipLaunchKernel(hipLaunchKernelStructFunc6, dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h6, result_d6); + hipLaunchKernel(hipLaunchKernelStructFunc7, dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h7, result_d7); + hipLaunchKernel(hipLaunchKernelStructFunc8, dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h8, result_d8); - // Validation part of the struct - hipMemcpy(hipLaunchKernelStruct_h1, hipLaunchKernelStruct_d1, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t1), hipMemcpyDeviceToHost); - for (int k = 0; k < STRUCT_SIZE; ++k) - HIPASSERT(hipLaunchKernelStruct_h1[k].result == true); + // Validation part of the struct, hipLaunchKernelStructFunc1 + hipMemcpy(result_h1, result_d1, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) + HIPASSERT(result_h1[k] == true); - // Validation part of the struct - hipMemcpy(hipLaunchKernelStruct_h2, hipLaunchKernelStruct_d2, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t2), hipMemcpyDeviceToHost); - for (int k = 0; k < STRUCT_SIZE; ++k) - HIPASSERT(hipLaunchKernelStruct_h2[k].result == true); + // Validation part of the struct, hipLaunchKernelStructFunc2 + hipMemcpy(result_h2, result_d2, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) + HIPASSERT(result_h2[k] == true); - // Validation part of the struct - hipMemcpy(hipLaunchKernelStruct_h3, hipLaunchKernelStruct_d3, STRUCT_SIZE*sizeof(hipLaunchKernelStruct_t3), hipMemcpyDeviceToHost); - for (int k = 0; k < STRUCT_SIZE; ++k) - HIPASSERT(hipLaunchKernelStruct_h3[k].result == true); + // Validation part of the struct, hipLaunchKernelStructFunc3 + hipMemcpy(result_h3, result_d3, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) + HIPASSERT(result_h3[k] == true); + + // Validation part of the struct, hipLaunchKernelStructFunc4 + hipMemcpy(result_h4, result_d4, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) + HIPASSERT(result_h4[k] == true); + + // Validation part of the struct, hipLaunchKernelStructFunc5 + hipMemcpy(result_h5, result_d5, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) + HIPASSERT(result_h5[k] == true); + + // 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) + HIPASSERT(result_h6[k] == true); + + // Validation part of the struct, hipLaunchKernelStructFunc7 + hipMemcpy(result_h7, result_d7, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) + HIPASSERT(result_h7[k] == true); + + // Validation part of the struct, hipLaunchKernelStructFunc7 + 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); // Test case with hipLaunchKernel inside another macro: float e0; @@ -196,5 +364,21 @@ int main() { MY_LAUNCH_WITH_PAREN(hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCall"); #endif + hipFree((void **)&result_h1); + hipFree((void **)&result_d1); + hipFree((void **)&result_h2); + hipFree((void **)&result_d2); + hipFree((void **)&result_h3); + hipFree((void **)&result_d3); + hipFree((void **)&result_h4); + hipFree((void **)&result_d4); + hipFree((void **)&result_h5); + hipFree((void **)&result_d5); + hipFree((void **)&result_h6); + hipFree((void **)&result_d6); + hipFree((void **)&result_h7); + hipFree((void **)&result_d7); + hipFree((void **)&result_h8); + hipFree((void **)&result_d8); passed(); } From a76d8f558096b5e841195fc92e64b627d1d65ceb Mon Sep 17 00:00:00 2001 From: Srinivasuluch Date: Wed, 27 Jun 2018 19:57:44 +0530 Subject: [PATCH 05/12] Update hipLaunchParm.cpp Added a macro to disable the alignment test as it's broken --- tests/src/kernel/hipLaunchParm.cpp | 20 ++++++++++++-------- 1 file changed, 12 insertions(+), 8 deletions(-) diff --git a/tests/src/kernel/hipLaunchParm.cpp b/tests/src/kernel/hipLaunchParm.cpp index 2b8a29c329..fcdcf6dc28 100644 --- a/tests/src/kernel/hipLaunchParm.cpp +++ b/tests/src/kernel/hipLaunchParm.cpp @@ -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; From a9419098dee11289b9af77c7d5e7e2d9ad6c707d Mon Sep 17 00:00:00 2001 From: Srinivasuluch Date: Wed, 27 Jun 2018 20:49:07 +0530 Subject: [PATCH 06/12] Update hipLaunchParm.cpp --- tests/src/kernel/hipLaunchParm.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/tests/src/kernel/hipLaunchParm.cpp b/tests/src/kernel/hipLaunchParm.cpp index fcdcf6dc28..5b53256b23 100644 --- a/tests/src/kernel/hipLaunchParm.cpp +++ b/tests/src/kernel/hipLaunchParm.cpp @@ -30,14 +30,14 @@ THE SOFTWARE. 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 +// This test is to verify Struct with variables to check the hipLaunchKernel() support, read from device. typedef struct hipLaunchKernelStruct1 { int li; // local int float lf; // local float bool result; // local bool } hipLaunchKernelStruct_t1; -// This test is to verify struct with padding, read and write into the same struct +// This test is to verify struct with padding, read from device typedef struct hipLaunchKernelStruct2 { char c1; long l1; @@ -46,7 +46,7 @@ typedef struct hipLaunchKernelStruct2 { bool result; } hipLaunchKernelStruct_t2; -// This test is to verify struct with padding, read and write into the same struct +// This test is to verify struct with padding, read from device typedef struct hipLaunchKernelStruct3 { char bf1; char bf2; From 9e78ef99e5533389433a8c4adfdafa39a656cadb Mon Sep 17 00:00:00 2001 From: Srinivasuluch Date: Wed, 27 Jun 2018 21:48:37 +0530 Subject: [PATCH 07/12] Update hipLaunchParm.cpp changing the condition, a = b, as Sam suggested --- tests/src/kernel/hipLaunchParm.cpp | 252 +++++++++++++++++++---------- 1 file changed, 164 insertions(+), 88 deletions(-) diff --git a/tests/src/kernel/hipLaunchParm.cpp b/tests/src/kernel/hipLaunchParm.cpp index 5b53256b23..1568d01dd0 100644 --- a/tests/src/kernel/hipLaunchParm.cpp +++ b/tests/src/kernel/hipLaunchParm.cpp @@ -30,7 +30,8 @@ THE SOFTWARE. static const int BLOCK_DIM_SIZE = 1024; -// This test is to verify Struct with variables to check the hipLaunchKernel() support, read from device. +// This test is to verify Struct with variables to check the hipLaunchKernel() +// support, read from device. typedef struct hipLaunchKernelStruct1 { int li; // local int float lf; // local float @@ -67,95 +68,137 @@ 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; +} __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 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 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; 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) { +// Passing struct to a hipLaunchKernel(), +// read and write into the same struct +__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 - result_d1[x] = ((hipLaunchKernelStruct_.li == 1) && (hipLaunchKernelStruct_.lf == 1.0) && (hipLaunchKernelStruct_.result == false)) ? true : false; + result_d1[x] = ((hipLaunchKernelStruct_.li == 1) + && (hipLaunchKernelStruct_.lf == 1.0) + && (hipLaunchKernelStruct_.result == false)); } -// 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) { +// 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) { int x = blockIdx.x * blockDim.x + threadIdx.x; // set the result to true if the condition met - result_d2[x] = ((hipLaunchKernelStruct_.c1 == 'a') && (hipLaunchKernelStruct_.l1 == 1.0) - && (hipLaunchKernelStruct_.c2 == 'b') && (hipLaunchKernelStruct_.l2 == 2.0) ) ? true : false; + result_d2[x] = ((hipLaunchKernelStruct_.c1 == 'a') + && (hipLaunchKernelStruct_.l1 == 1.0) + && (hipLaunchKernelStruct_.c2 == 'b') + && (hipLaunchKernelStruct_.l2 == 2.0) ); } -// 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) { +// 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) { int x = blockIdx.x * blockDim.x + threadIdx.x; // set the result to true if the condition met - result_d3[x] = ((hipLaunchKernelStruct_.bf1 == 1) && (hipLaunchKernelStruct_.bf2 == 1) - && (hipLaunchKernelStruct_.l1 == 1.0) && (hipLaunchKernelStruct_.bf3 == 1) ) ? true : false; + result_d3[x] = ((hipLaunchKernelStruct_.bf1 == 1) + && (hipLaunchKernelStruct_.bf2 == 1) + && (hipLaunchKernelStruct_.l1 == 1.0) + && (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 the 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 - result_d4[x] = (sizeof(hipLaunchKernelStruct_) == 1) ? true : false; + result_d4[x] = (sizeof(hipLaunchKernelStruct_) == 1); } // 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 - result_d5[x] = ((hipLaunchKernelStruct_.c1 == 'c') && (*hipLaunchKernelStruct_.cp == 'p')) ? true : false; + result_d5[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (*hipLaunchKernelStruct_.cp == 'p')); } -// 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) { +// 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) { int x = blockIdx.x * blockDim.x + threadIdx.x; // 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; + // get the address of the struct + // size_t(p)%8 will be 0 if aligned to 8Byte address space + int *p = (int*)(&hipLaunchKernelStruct_); + result_d6[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (hipLaunchKernelStruct_.si == 1) + && ((size_t(p))%8 ==0)); } -// Passing struct which is aligned to 16Byte to a hipLaunchKernel(), set the result_d7 to true if condition met -__global__ void hipLaunchKernelStructFunc7(hipLaunchParm lp, hipLaunchKernelStruct_t7 hipLaunchKernelStruct_, bool* result_d7) { +// Passing struct which is aligned to 16Byte to a hipLaunchKernel(), +// set the result_d7 to true if condition met +__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 - 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; + // get the address of the struct + // size_t(p)%16 will be 0 if aligned to 16Byte address space + int *p = (int*)(&hipLaunchKernelStruct_); + result_d7[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (hipLaunchKernelStruct_.si == 1) + && ((size_t(p))%16 ==0) ); } -// Passing struct which is packed & aligned to 4Byte to a hipLaunchKernel(), set the result_d8 to true if condition met -__global__ void hipLaunchKernelStructFunc8(hipLaunchParm lp, hipLaunchKernelStruct_t8 hipLaunchKernelStruct_, bool* result_d8) { +// Passing struct which is packed & aligned to 4Byte to a hipLaunchKernel(), +// set the result_d8 to true if condition met +__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 - int *p = (int*)(&hipLaunchKernelStruct_); // get the address of the xth element, struct[x], - // 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; + // get the address of the xth element, struct[x], + // size_t(p)%4 will be 0 if aligned to 4Byte address space + int *p = (int*)(&hipLaunchKernelStruct_); + result_d8[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (hipLaunchKernelStruct_.si == 1) + && ((size_t(p))%4 ==0) ); } __global__ void vAdd(hipLaunchParm lp, float* a) {} @@ -165,29 +208,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); \ } @@ -202,7 +245,8 @@ 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 is accessible from device. } hipLaunchKernelStruct_h1.li = 1; hipLaunchKernelStruct_h1.lf = 1.0; @@ -214,7 +258,8 @@ int main() { hipMalloc((void**)&result_d2, BLOCK_DIM_SIZE*sizeof(bool)); hipHostMalloc((void**)&result_h2, 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 is accessible from device. } hipLaunchKernelStruct_h2.c1 = 'a'; hipLaunchKernelStruct_h2.l1 = 1.0; @@ -228,13 +273,15 @@ 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 is accessible from device. } 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; // This will be set to true if + // the the condition is satisfied, from device side // empty struct hipLaunchKernelStruct_t4 hipLaunchKernelStruct_h4; @@ -242,21 +289,25 @@ 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() hipLaunchKernelStruct_t5 hipLaunchKernelStruct_h5; - char* cp_d5; // This is passed as pointer to struct member, struct.cp = &cp_d5 + // This is passed as pointer to struct member, struct.cp = &cp_d5 + char* cp_d5; bool *result_d5, *result_h5; hipMalloc((void**)&result_d5, BLOCK_DIM_SIZE*sizeof(bool)); - hipMalloc((void**)&cp_d5, sizeof(char)); // allocating memory for char pointer on device + // allocating memory for char pointer on device + hipMalloc((void**)&cp_d5, sizeof(char)); hipHostMalloc((void**)&result_h5, BLOCK_DIM_SIZE*sizeof(bool)); *cp_d5 = 'p'; // initializing memory to 'p' 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) @@ -267,7 +318,8 @@ 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) @@ -278,7 +330,8 @@ 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; @@ -288,7 +341,8 @@ 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: @@ -296,68 +350,89 @@ int main() { hipLaunchKernel(vAdd, 1024, dim3(1), 0, 0, Ad); 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); - hipLaunchKernel(hipLaunchKernelStructFunc2, dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h2, result_d2); - hipLaunchKernel(hipLaunchKernelStructFunc3, dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h3, result_d3); - hipLaunchKernel(hipLaunchKernelStructFunc4, dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h4, result_d4); - hipLaunchKernel(hipLaunchKernelStructFunc5, dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h5, result_d5); - hipLaunchKernel(hipLaunchKernelStructFunc6, dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h6, result_d6); - hipLaunchKernel(hipLaunchKernelStructFunc7, dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h7, result_d7); - hipLaunchKernel(hipLaunchKernelStructFunc8, dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h8, result_d8); + hipLaunchKernel(hipLaunchKernelStructFunc1, dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h1, result_d1); + hipLaunchKernel(hipLaunchKernelStructFunc2, dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h2, result_d2); + hipLaunchKernel(hipLaunchKernelStructFunc3, dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h3, result_d3); + hipLaunchKernel(hipLaunchKernelStructFunc4, dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h4, result_d4); + hipLaunchKernel(hipLaunchKernelStructFunc5, dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h5, result_d5); + hipLaunchKernel(hipLaunchKernelStructFunc6, dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h6, result_d6); + hipLaunchKernel(hipLaunchKernelStructFunc7, dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h7, result_d7); + hipLaunchKernel(hipLaunchKernelStructFunc8, dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h8, result_d8); // Validation part of the struct, hipLaunchKernelStructFunc1 - hipMemcpy(result_h1, result_d1, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost); + hipMemcpy(result_h1, result_d1, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyDeviceToHost); for (int k = 0; k < BLOCK_DIM_SIZE; ++k) HIPASSERT(result_h1[k] == true); // Validation part of the struct, hipLaunchKernelStructFunc2 - hipMemcpy(result_h2, result_d2, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost); + hipMemcpy(result_h2, result_d2, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyDeviceToHost); for (int k = 0; k < BLOCK_DIM_SIZE; ++k) HIPASSERT(result_h2[k] == true); // Validation part of the struct, hipLaunchKernelStructFunc3 - hipMemcpy(result_h3, result_d3, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost); + hipMemcpy(result_h3, result_d3, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyDeviceToHost); for (int k = 0; k < BLOCK_DIM_SIZE; ++k) HIPASSERT(result_h3[k] == true); // Validation part of the struct, hipLaunchKernelStructFunc4 - hipMemcpy(result_h4, result_d4, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost); + hipMemcpy(result_h4, result_d4, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyDeviceToHost); for (int k = 0; k < BLOCK_DIM_SIZE; ++k) HIPASSERT(result_h4[k] == true); // Validation part of the struct, hipLaunchKernelStructFunc5 - hipMemcpy(result_h5, result_d5, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost); + hipMemcpy(result_h5, result_d5, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyDeviceToHost); 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 + // alignment is broken hence disabled the validation part + #if DISABLE_TEST // Validation part of the struct, hipLaunchKernelStructFunc6 - hipMemcpy(result_h6, result_d6, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost); + hipMemcpy(result_h6, result_d6, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyDeviceToHost); for (int k = 0; k < BLOCK_DIM_SIZE; ++k) HIPASSERT(result_h6[k] == true); // Validation part of the struct, hipLaunchKernelStructFunc7 - hipMemcpy(result_h7, result_d7, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost); + hipMemcpy(result_h7, result_d7, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyDeviceToHost); for (int k = 0; k < BLOCK_DIM_SIZE; ++k) HIPASSERT(result_h7[k] == true); // Validation part of the struct, hipLaunchKernelStructFunc7 - hipMemcpy(result_h8, result_d8, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost); + 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; - GPU_PRINT_TIME(hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), e0, j); - GPU_PRINT_TIME(WRAP(hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad)), e0, j); + GPU_PRINT_TIME(hipLaunchKernel(vAdd, dim3(1024), + dim3(1), 0, 0, Ad), e0, j); + GPU_PRINT_TIME(WRAP(hipLaunchKernel(vAdd, dim3(1024), + dim3(1), 0, 0, Ad)), e0, j); #ifdef EXTRA_PARENS_1 // Don't wrap hipLaunchKernel in extra set of parens: - GPU_PRINT_TIME((hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad)), e0, j); + GPU_PRINT_TIME((hipLaunchKernel(vAdd, dim3(1024), + dim3(1), 0, 0, Ad)), e0, j); #endif - MY_LAUNCH(hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCall"); + MY_LAUNCH(hipLaunchKernel(vAdd, dim3(1024), dim3(1), + 0, 0, Ad), true, "firstCall"); float* A; float e1; @@ -365,7 +440,8 @@ int main() { #ifdef EXTRA_PARENS_2 // MY_LAUNCH_WITH_PAREN wraps cmd in () which can cause issues. - MY_LAUNCH_WITH_PAREN(hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCall"); + MY_LAUNCH_WITH_PAREN(hipLaunchKernel(vAdd, dim3(1024), + dim3(1), 0, 0, Ad), true, "firstCall"); #endif hipFree((void **)&result_h1); From 8b67625b65cbaa73317bc0e87a57763dd52655cb Mon Sep 17 00:00:00 2001 From: Srinivasuluch Date: Wed, 27 Jun 2018 22:04:12 +0530 Subject: [PATCH 08/12] Update hipLaunchParm.cpp wrap up to 80 characters per line --- tests/src/kernel/hipLaunchParm.cpp | 178 ++++++++++++++++------------- 1 file changed, 101 insertions(+), 77 deletions(-) diff --git a/tests/src/kernel/hipLaunchParm.cpp b/tests/src/kernel/hipLaunchParm.cpp index 1568d01dd0..8df10806cb 100644 --- a/tests/src/kernel/hipLaunchParm.cpp +++ b/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), From 181ccebee6dc8b2da0352e9c427c9f7f9470c2d9 Mon Sep 17 00:00:00 2001 From: Srinivasuluch Date: Wed, 27 Jun 2018 22:06:53 +0530 Subject: [PATCH 09/12] Update hipLaunchParm.cpp Changing macro name --- tests/src/kernel/hipLaunchParm.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tests/src/kernel/hipLaunchParm.cpp b/tests/src/kernel/hipLaunchParm.cpp index 8df10806cb..44c31b0098 100644 --- a/tests/src/kernel/hipLaunchParm.cpp +++ b/tests/src/kernel/hipLaunchParm.cpp @@ -26,7 +26,7 @@ THE SOFTWARE. #include "hip/hip_runtime.h" #include "test_common.h" -#define DISABLE_TEST 0 +#define ENABLE_ALIGNMENT_TEST 0 static const int BLOCK_DIM_SIZE = 1024; @@ -422,7 +422,7 @@ int main() { HIPASSERT(result_h5[k] == true); // alignment is broken hence disabled the validation part - #if DISABLE_TEST + #if ENABLE_ALIGNMENT_TEST // Validation part of the struct, hipLaunchKernelStructFunc6 hipMemcpy(result_h6, result_d6, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost); From 19ace627a3e911a5c0ffb620b9972e140fecc89b Mon Sep 17 00:00:00 2001 From: srinivas Charupally Date: Thu, 12 Jul 2018 17:08:41 +0530 Subject: [PATCH 10/12] Adding more struct scenarions and a ResultValidation() --- tests/src/kernel/hipLaunchParm.cpp | 702 +++++++++++++++++++++-------- 1 file changed, 502 insertions(+), 200 deletions(-) diff --git a/tests/src/kernel/hipLaunchParm.cpp b/tests/src/kernel/hipLaunchParm.cpp index 44c31b0098..420a74e7fa 100644 --- a/tests/src/kernel/hipLaunchParm.cpp +++ b/tests/src/kernel/hipLaunchParm.cpp @@ -23,13 +23,48 @@ THE SOFTWARE. * HIT_END */ +#include #include "hip/hip_runtime.h" #include "test_common.h" +// Memory alignment is broken #define ENABLE_ALIGNMENT_TEST 0 +// Accessing struct class object from device is broken +#define ENABLE_CLASS_OBJ_ACCESS 0 +// accessing dynamic/heap memory from device is broken +#define ENABLE_HEAP_MEMORY_ACCESS 0 +// STL implementation broken, accessing dynamic/heap memory fail +// broken on hcc, working fine on hip-clang +#define ENABLE_USER_STL 0 +// out of order initalization is broken +#define ENABLE_OUT_OF_ORDER_INITIALIZATION 0 +// Direct initialization of struct broken, +// ip_d9 is a pointer, uint_t*, hipLaunchKernelStruct_h9 = {'c', ip_d9}; +#define ENABLE_DECLARE_INITIALIZATION_POINTER 0 static const int BLOCK_DIM_SIZE = 1024; +// allocate memory on device and host for result validation +static bool *result_d, *result_h; +static hipError_t hipMallocError = hipMalloc((void**)&result_d, + BLOCK_DIM_SIZE*sizeof(bool)); +static hipError_t hipHostMallocError = hipHostMalloc((void**)&result_h, + BLOCK_DIM_SIZE*sizeof(bool)); +static hipError_t hipMemsetError = hipMemset(result_d, + false, BLOCK_DIM_SIZE); + +static void ResultValidation() { + hipMemcpy(result_h, result_d, BLOCK_DIM_SIZE*sizeof(bool), + hipMemcpyDeviceToHost); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) + HIPASSERT(result_h[k] == true); + + // reset the memory to false to reuse it. + hipMemset(result_d, false, BLOCK_DIM_SIZE); + hipMemset(result_h, false, BLOCK_DIM_SIZE); + return; +} + // This test is to verify Struct with variables // support, read from device. typedef struct hipLaunchKernelStruct1 { @@ -90,10 +125,120 @@ typedef struct hipLaunchKernelStruct8 { bool b; }__attribute__((packed, aligned(4))) hipLaunchKernelStruct_t8; -// Passing struct to a hipLaunchKernel(), +// This test is to verify const struct object +typedef struct hipLaunchKernelStruct9 { + char c1; + uint32_t* ip; // uint pointer +} hipLaunchKernelStruct_t9; + +// This test is to verify struct with stdint types, uintN_t +typedef struct hipLaunchKernelStruct10 { + uint64_t u64; + uint32_t u32; + uint8_t u8; +} hipLaunchKernelStruct_t10; + +// This test is to verify struct with volatile member +typedef struct hipLaunchKernelStruct11 { + int i1; + volatile unsigned int vint; +} hipLaunchKernelStruct_t11; + +// This test is to verify struct with simple class object +class base { + public: + int i = 0; + base() {} +}; +typedef struct hipLaunchKernelStruct12 { + base b; + char c1; +} hipLaunchKernelStruct_t12; + +// This test is to verify struct with __device__ func() attribute +typedef struct hipLaunchKernelStruct13 { + int i1; + __device__ int getvalue() { return i1; } +} hipLaunchKernelStruct_t13; + +// This test is to verify struct with array variable, +// write to from device +typedef struct hipLaunchKernelStruct14 { + int readint; + int writeint[BLOCK_DIM_SIZE]; // will write to this from device +} hipLaunchKernelStruct_t14; + +// This test is to verify struct with dynamic memory, new int +// the heap memory will be accessed from device +typedef struct hipLaunchKernelStruct15 { + char c1; + int* heapmem = new int[BLOCK_DIM_SIZE]; +} hipLaunchKernelStruct_t15; + +// This test is to verify simple template struct +template +struct hipLaunchKernelStruct_t16 { + T t1; +}; + +// This test is to verify simple explicity template struct +template struct hipLaunchKernelStruct_t17 {}; +template<> // explicit template +struct hipLaunchKernelStruct_t17 { + int t1; +}; + +// This test is to verity write to struct memory using __device__ func() +typedef struct hipLaunchKernelStruct18 { + char c1; + __device__ void setChar(char c) { c1 = c; } + __device__ int getChar() { return c1; } +} hipLaunchKernelStruct_t18; + +// This test is to verity user defined STL, simple stack implementation +typedef struct stackNode { + int data; + stackNode* nextNode = NULL; +} stackNode_t; +typedef struct hipLaunchKernelStruct19 { + stackNode_t* stack = NULL; + unsigned int size_ = 0; + void pushMe(int value) { // not a device function, setting from host + stackNode_t* newNode = new stackNode_t; + newNode->data = value; + ++size_; + if (stack == NULL) { + stack = newNode; + return; + } + stackNode_t* currentHead = stack; + stack = newNode; + stack->nextNode = currentHead; + return; + } + __device__ void popMe() { + stackNode_t* currentHead = stack; + stack = stack->nextNode; + --size_; + // delete currentHead; // no idea why delete not working + return; + } + int stackSize() { + return size_; + } +} hipLaunchKernelStruct_t19; + +// This test is to verify out of order initalizer of struct elements +// and access in-order, from device. +typedef struct hipLaunchKernelStruct20 { + char name; + int age; + int rank; +} hipLaunchKernelStruct_t20; + +// Passing struct to a hipLaunchKernelGGL(), // read and write into the same struct __global__ void hipLaunchKernelStructFunc1( - hipLaunchParm lp, hipLaunchKernelStruct_t1 hipLaunchKernelStruct_, bool* result_d1) { int x = blockIdx.x * blockDim.x + threadIdx.x; @@ -104,10 +249,9 @@ __global__ void hipLaunchKernelStructFunc1( && (hipLaunchKernelStruct_.result == false)); } -// Passing struct to a hipLaunchKernel(), checks padding, +// Passing struct to a hipLaunchKernelGGL(), checks padding, // read and write into the same struct __global__ void hipLaunchKernelStructFunc2( - hipLaunchParm lp, hipLaunchKernelStruct_t2 hipLaunchKernelStruct_, bool* result_d2) { int x = blockIdx.x * blockDim.x + threadIdx.x; @@ -119,10 +263,9 @@ __global__ void hipLaunchKernelStructFunc2( && (hipLaunchKernelStruct_.l2 == 2.0) ); } -// Passing struct to a hipLaunchKernel(), checks padding, +// Passing struct to a hipLaunchKernelGGL(), checks padding, // read and write into the same struct __global__ void hipLaunchKernelStructFunc3( - hipLaunchParm lp, hipLaunchKernelStruct_t3 hipLaunchKernelStruct_, bool* result_d3) { int x = blockIdx.x * blockDim.x + threadIdx.x; @@ -134,10 +277,9 @@ __global__ void hipLaunchKernelStructFunc3( && (hipLaunchKernelStruct_.bf3 == 1) ); } -// Passing empty struct to a hipLaunchKernel(), +// Passing empty struct to a hipLaunchKernelGGL(), // 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; @@ -146,9 +288,8 @@ __global__ void hipLaunchKernelStructFunc4( result_d4[x] = (sizeof(hipLaunchKernelStruct_) == 1); } -// Passing struct with pointer object to a hipLaunchKernel() +// Passing struct with pointer object to a hipLaunchKernelGGL() __global__ void hipLaunchKernelStructFunc5( - hipLaunchParm lp, hipLaunchKernelStruct_t5 hipLaunchKernelStruct_, bool* result_d5) { int x = blockIdx.x * blockDim.x + threadIdx.x; @@ -158,10 +299,9 @@ __global__ void hipLaunchKernelStructFunc5( && (*hipLaunchKernelStruct_.cp == 'p')); } -// Passing struct which is aligned to 8Byte to a hipLaunchKernel(), +// Passing struct which is aligned to 8Byte to a hipLaunchKernelGGL(), // set the result_d6 to true if condition met __global__ void hipLaunchKernelStructFunc6( - hipLaunchParm lp, hipLaunchKernelStruct_t6 hipLaunchKernelStruct_, bool* result_d6) { int x = blockIdx.x * blockDim.x + threadIdx.x; @@ -178,7 +318,6 @@ __global__ void hipLaunchKernelStructFunc6( // 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) { int x = blockIdx.x * blockDim.x + threadIdx.x; @@ -195,7 +334,6 @@ __global__ void hipLaunchKernelStructFunc7( // 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) { int x = blockIdx.x * blockDim.x + threadIdx.x; @@ -209,7 +347,140 @@ __global__ void hipLaunchKernelStructFunc8( && ((size_t(p))%4 ==0) ); } -__global__ void vAdd(hipLaunchParm lp, float* a) {} +// Passing struct with uint pointer object to a hipLaunchKernelGGL() +__global__ void hipLaunchKernelStructFunc9( + const hipLaunchKernelStruct_t9 hipLaunchKernelStruct_, + bool* result_d9) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + result_d9[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (*hipLaunchKernelStruct_.ip == 1)); +} + +// Passing struct with stdint types object, uintN_t, to a hipLaunchKernelGGL() +__global__ void hipLaunchKernelStructFunc10( + hipLaunchKernelStruct_t10 hipLaunchKernelStruct_, + bool* result_d10) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + result_d10[x] = ((hipLaunchKernelStruct_.u64 == UINT64_MAX) + && (hipLaunchKernelStruct_.u32 == 1) + && (hipLaunchKernelStruct_.u8 == UINT8_MAX)); +} + +// Passing struct with volatile member, to a hipLaunchKernelGGL() +__global__ void hipLaunchKernelStructFunc11( + hipLaunchKernelStruct_t11 hipLaunchKernelStruct_, + bool* result_d11) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + result_d11[x] = ((hipLaunchKernelStruct_.i1 == 1) + && (hipLaunchKernelStruct_.vint == 0)); +} + +// Passing struct with simple class obj, to a hipLaunchKernelGGL() +__global__ void hipLaunchKernelStructFunc12( + hipLaunchKernelStruct_t12 hipLaunchKernelStruct_, + bool* result_d12) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + result_d12[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (hipLaunchKernelStruct_.b.i == 0)); +} + +// Passing struct with simple __device__ func(), to a hipLaunchKernelGGL() +__global__ void hipLaunchKernelStructFunc13( + hipLaunchKernelStruct_t13 hipLaunchKernelStruct_, + bool* result_d13) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + result_d13[x] = ((hipLaunchKernelStruct_.i1 == 1) + && (hipLaunchKernelStruct_.getvalue() == 1)); +} + +// Passing struct with array variable, write to from device +__global__ void hipLaunchKernelStructFunc14( + hipLaunchKernelStruct_t14 hipLaunchKernelStruct_, + bool* result_d14) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + hipLaunchKernelStruct_.writeint[x] = 1; + // set the result to true if the condition met + result_d14[x] = ((hipLaunchKernelStruct_.readint == 1) + && (hipLaunchKernelStruct_.writeint[x] == 1)); +} + +// Passing struct with struct with dynamic memory, new int +// the heap memory will be accessed from device +__global__ void hipLaunchKernelStructFunc15( + hipLaunchKernelStruct_t15 hipLaunchKernelStruct_, + bool* result_d15) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + result_d15[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (hipLaunchKernelStruct_.heapmem[x] == 1)); +} + +// Passing simple template struct +__global__ void hipLaunchKernelStructFunc16( + hipLaunchKernelStruct_t16 hipLaunchKernelStruct_, + bool* result_d16) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + result_d16[x] = (hipLaunchKernelStruct_.t1 == 'c'); +} + +// Passing simple explicit template struct +__global__ void hipLaunchKernelStructFunc17( + hipLaunchKernelStruct_t17 hipLaunchKernelStruct_, + bool* result_d17) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + result_d17[x] = (hipLaunchKernelStruct_.t1 == 1); +} + +// Passing struct and write to struct memory using __device__ func() +__global__ void hipLaunchKernelStructFunc18( + hipLaunchKernelStruct_t18 hipLaunchKernelStruct_, + bool* result_d18) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + hipLaunchKernelStruct_.setChar('c'); + // set the result to true if the condition met + result_d18[x] = (hipLaunchKernelStruct_.getChar() == 'c'); +} + +// Passing simple user defined stack implemenration, using __device__ func() +__global__ void hipLaunchKernelStructFunc19( + hipLaunchKernelStruct_t19 hipLaunchKernelStruct_) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // stack should be empty after the kernel execustion, verify on host side + hipLaunchKernelStruct_.popMe(); +} + +// Passing out of order initalized struct, access in-order +__global__ void hipLaunchKernelStructFunc20( + hipLaunchKernelStruct_t20 hipLaunchKernelStruct_, + bool* result_d20) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // accessing struct members in order + result_d20[x] = (hipLaunchKernelStruct_.name = 'A' + // strcmp(hipLaunchKernelStruct_.name, "AMD") -> strcmp is not broken + && hipLaunchKernelStruct_.age == 42 + && hipLaunchKernelStruct_.rank == 2); +} + +__global__ void vAdd(float* a) {} //--- // Some wrapper macro for testing: @@ -243,219 +514,263 @@ __global__ void vAdd(hipLaunchParm lp, float* a) {} int main() { - float* Ad; + // Validating memory & initial value, for result_d, result_h + HIPASSERT(hipMallocError == hipSuccess); + HIPASSERT(hipHostMallocError == hipSuccess); + HIPASSERT(hipMemsetError == hipSuccess); - hipMalloc((void**)&Ad, 1024); - - // Struct type, check access from device. + // Test: Passing Struct type, check access from device. hipLaunchKernelStruct_t1 hipLaunchKernelStruct_h1; - bool *result_d1, *result_h1; - 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 size is 1Byte, from device size - } hipLaunchKernelStruct_h1.li = 1; hipLaunchKernelStruct_h1.lf = 1.0; hipLaunchKernelStruct_h1.result = false; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc1), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h1, + result_d); + ResultValidation(); - // Struct type, checks padding + // Test: Passing Struct type, checks padding hipLaunchKernelStruct_t2 hipLaunchKernelStruct_h2; - bool *result_d2, *result_h2; - hipMalloc((void**)&result_d2, BLOCK_DIM_SIZE*sizeof(bool)); - hipHostMalloc((void**)&result_h2, 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. - } hipLaunchKernelStruct_h2.c1 = 'a'; hipLaunchKernelStruct_h2.l1 = 1.0; hipLaunchKernelStruct_h2.c2 = 'b'; hipLaunchKernelStruct_h2.l2 = 2.0; hipLaunchKernelStruct_h2.result = false; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc2), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h2, + result_d); + ResultValidation(); - // Struct type, checks padding, assigning integer to a char + // Test: Passing Struct type, checks padding, assigning integer to a char hipLaunchKernelStruct_t3 hipLaunchKernelStruct_h3; - bool *result_d3, *result_h3; - 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 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; + hipLaunchKernelStruct_h3.result = false; // initialize to false, will be set to // true if the struct size is 1Byte, from device size + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc3), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h3, + result_d); + ResultValidation(); - // empty struct + // Test: Passing empty struct hipLaunchKernelStruct_t4 hipLaunchKernelStruct_h4; - bool *result_d4, *result_h4; - 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 - } + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc4), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h4, + result_d); + ResultValidation(); - // Passing struct with pointer object to a hipLaunchKernel() + // Test: Passing struct with pointer object to a hipLaunchKernelGGL() hipLaunchKernelStruct_t5 hipLaunchKernelStruct_h5; - // This is passed as pointer to struct member, struct.cp = &cp_d5 - char* cp_d5; - bool *result_d5, *result_h5; - hipMalloc((void**)&result_d5, BLOCK_DIM_SIZE*sizeof(bool)); + char* cp_d5; // This is passed as pointer to struct member // allocating memory for char pointer on device - hipMalloc((void**)&cp_d5, sizeof(char)); - hipHostMalloc((void**)&result_h5, BLOCK_DIM_SIZE*sizeof(bool)); - *cp_d5 = 'p'; // initializing memory to 'p' + HIPCHECK(hipMalloc((void**)&cp_d5, sizeof(char))); + HIPCHECK(hipMemset(cp_d5, 'p', sizeof(char))); 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 - } + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc5), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h5, + result_d); + ResultValidation(); - // Passing struct with aligned(8) + // Test: Passing struct with aligned(8) hipLaunchKernelStruct_t6 hipLaunchKernelStruct_h6; - bool *result_d6, *result_h6; - hipMalloc((void**)&result_d6, BLOCK_DIM_SIZE*sizeof(bool)); - hipHostMalloc((void**)&result_h6, BLOCK_DIM_SIZE*sizeof(bool)); 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 - } - - // Passing struct with aligned(16) - hipLaunchKernelStruct_t7 hipLaunchKernelStruct_h7; - bool *result_d7, *result_h7; - hipMalloc((void**)&result_d7, BLOCK_DIM_SIZE*sizeof(bool)); - hipHostMalloc((void**)&result_h7, BLOCK_DIM_SIZE*sizeof(bool)); - 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 - } - // Passing struct with packed aligned to 6Bytes - hipLaunchKernelStruct_t8 hipLaunchKernelStruct_h8; - bool *result_d8, *result_h8; - hipMalloc((void**)&result_d8, BLOCK_DIM_SIZE*sizeof(bool)); - hipHostMalloc((void**)&result_h8, BLOCK_DIM_SIZE*sizeof(bool)); - 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 - } - - // Test the different hipLaunchParm options: - hipLaunchKernel(vAdd, size_t(1024), 1, 0, 0, Ad); - hipLaunchKernel(vAdd, 1024, dim3(1), 0, 0, Ad); - 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); - hipLaunchKernel(hipLaunchKernelStructFunc2, dim3(BLOCK_DIM_SIZE), - dim3(1), 0, 0, hipLaunchKernelStruct_h2, - result_d2); - hipLaunchKernel(hipLaunchKernelStructFunc3, dim3(BLOCK_DIM_SIZE), - dim3(1), 0, 0, hipLaunchKernelStruct_h3, - result_d3); - hipLaunchKernel(hipLaunchKernelStructFunc4, dim3(BLOCK_DIM_SIZE), - dim3(1), 0, 0, hipLaunchKernelStruct_h4, - result_d4); - hipLaunchKernel(hipLaunchKernelStructFunc5, dim3(BLOCK_DIM_SIZE), - dim3(1), 0, 0, hipLaunchKernelStruct_h5, - result_d5); - hipLaunchKernel(hipLaunchKernelStructFunc6, dim3(BLOCK_DIM_SIZE), + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc6), + dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h6, - result_d6); - hipLaunchKernel(hipLaunchKernelStructFunc7, dim3(BLOCK_DIM_SIZE), - dim3(1), 0, 0, hipLaunchKernelStruct_h7, - result_d7); - hipLaunchKernel(hipLaunchKernelStructFunc8, dim3(BLOCK_DIM_SIZE), - dim3(1), 0, 0, hipLaunchKernelStruct_h8, - result_d8); - - // Validation part of the struct, hipLaunchKernelStructFunc1 - hipMemcpy(result_h1, result_d1, BLOCK_DIM_SIZE*sizeof(bool), - hipMemcpyDeviceToHost); - for (int k = 0; k < BLOCK_DIM_SIZE; ++k) - HIPASSERT(result_h1[k] == true); - - // Validation part of the struct, hipLaunchKernelStructFunc2 - hipMemcpy(result_h2, result_d2, BLOCK_DIM_SIZE*sizeof(bool), - hipMemcpyDeviceToHost); - for (int k = 0; k < BLOCK_DIM_SIZE; ++k) - HIPASSERT(result_h2[k] == true); - - // Validation part of the struct, hipLaunchKernelStructFunc3 - hipMemcpy(result_h3, result_d3, BLOCK_DIM_SIZE*sizeof(bool), - hipMemcpyDeviceToHost); - for (int k = 0; k < BLOCK_DIM_SIZE; ++k) - HIPASSERT(result_h3[k] == true); - - // Validation part of the struct, hipLaunchKernelStructFunc4 - hipMemcpy(result_h4, result_d4, BLOCK_DIM_SIZE*sizeof(bool), - hipMemcpyDeviceToHost); - for (int k = 0; k < BLOCK_DIM_SIZE; ++k) - HIPASSERT(result_h4[k] == true); - - // Validation part of the struct, hipLaunchKernelStructFunc5 - hipMemcpy(result_h5, result_d5, BLOCK_DIM_SIZE*sizeof(bool), - hipMemcpyDeviceToHost); - for (int k = 0; k < BLOCK_DIM_SIZE; ++k) - HIPASSERT(result_h5[k] == true); - + result_d); // alignment is broken hence disabled the validation part #if ENABLE_ALIGNMENT_TEST - // 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) - HIPASSERT(result_h6[k] == true); + ResultValidation(); + #endif - // Validation part of the struct, hipLaunchKernelStructFunc7 - hipMemcpy(result_h7, result_d7, BLOCK_DIM_SIZE*sizeof(bool), - hipMemcpyDeviceToHost); - for (int k = 0; k < BLOCK_DIM_SIZE; ++k) - HIPASSERT(result_h7[k] == true); - // Validation part of the struct, hipLaunchKernelStructFunc7 - 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: Passing struct with aligned(16) + hipLaunchKernelStruct_t7 hipLaunchKernelStruct_h7; + hipLaunchKernelStruct_h7.c1 = 'c'; + hipLaunchKernelStruct_h7.si = 1; + #if ENABLE_ALIGNMENT_TEST // This is broken on small bar + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc7), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h7, + result_d); + ResultValidation(); + #endif - // Test case with hipLaunchKernel inside another macro: + // Test: Passing struct with packed aligned to 6Bytes + hipLaunchKernelStruct_t8 hipLaunchKernelStruct_h8; + hipLaunchKernelStruct_h8.c1 = 'c'; + hipLaunchKernelStruct_h8.si = 1; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc8), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h8, + result_d); + // alignment is broken hence disabled the validation part + #if ENABLE_ALIGNMENT_TEST + ResultValidation(); + #endif + + + // Test: Passing const struct object to a hipLaunchKernelGGL() + uint32_t* ip_d9; + // allocating memory for char pointer on device + HIPCHECK(hipMalloc((void**)&ip_d9, sizeof(uint32_t))); + HIPCHECK(hipMemset(ip_d9, 1, sizeof(uint32_t))); + // ip_d9 passed as pointer to struct member, struct.ip = &ip_d9 + const hipLaunchKernelStruct_t9 hipLaunchKernelStruct_h9 = {'c', ip_d9}; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc9), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h9, + result_d); + #if ENABLE_DECLARE_INITIALIZATION_POINTER + ResultValidation(); + #endif + + + // Test: Passing struct with uintN_t as member variables + hipLaunchKernelStruct_t10 hipLaunchKernelStruct_h10; + hipLaunchKernelStruct_h10.u64 = UINT64_MAX; + hipLaunchKernelStruct_h10.u32 = 1; + hipLaunchKernelStruct_h10.u8 = UINT8_MAX; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc10), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h10, + result_d); + ResultValidation(); + + + // Test: Passing struct with uintN_t as member variables + hipLaunchKernelStruct_t11 hipLaunchKernelStruct_h11; + hipLaunchKernelStruct_h11.i1 = 1; + hipLaunchKernelStruct_h11.vint = 0; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc11), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h11, + result_d); + ResultValidation(); + + // Test: Passing struct with simple class object + hipLaunchKernelStruct_t12 hipLaunchKernelStruct_h12; + hipLaunchKernelStruct_h12.c1 = 'c'; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc12), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h12, + result_d); + #if ENABLE_CLASS_OBJ_ACCESS // access class obj from device broken + // Validation part of the struct, hipLaunchKernelStructFunc12 + ResultValidation(); + #endif + + // Test: Passing struct with simple __device__ func() + hipLaunchKernelStruct_t13 hipLaunchKernelStruct_h13; + hipLaunchKernelStruct_h13.i1 = 1; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc13), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h13, + result_d); + ResultValidation(); + + // Test: Passing struct with array variable, write to from device + hipLaunchKernelStruct_t14 hipLaunchKernelStruct_h14; + hipLaunchKernelStruct_h14.readint = 1; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc14), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h14, + result_d); + ResultValidation(); + + + // Test: Passing struct with heap memory, read to from device + hipLaunchKernelStruct_t15 hipLaunchKernelStruct_h15; + hipLaunchKernelStruct_h15.c1 = 'c'; + #if ENABLE_HEAP_MEMORY_ACCESS // causing page fault here,validation failed + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc15), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h15, + result_d); + ResultValidation(); + #endif + + // Test: Passing simple template struct + hipLaunchKernelStruct_t16 hipLaunchKernelStruct_h16; + hipLaunchKernelStruct_h16.t1 = 'c'; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc16), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h16, + result_d); + ResultValidation(); + + // Test: Passing simple explicit template struct + hipLaunchKernelStruct_t17 hipLaunchKernelStruct_h17; + hipLaunchKernelStruct_h17.t1 = 1; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc17), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h17, + result_d); + ResultValidation(); + + // Test: Passing struct with simple __device__ func() to struct memory + hipLaunchKernelStruct_t18 hipLaunchKernelStruct_h18; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc18), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h18, + result_d); + ResultValidation(); + + // Test: Passing user defined stack, + hipLaunchKernelStruct_t19 hipLaunchKernelStruct_h19; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc19), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h19); + #if ENABLE_USER_STL + // Validation part of the struct, hipLaunchKernelStructFunc19 + HIPASSERT(hipLaunchKernelStruct_h19.stackSize() == 0); + #endif + + // Test: Passing struct which is initiazed out of order + // accessing same elements in order from device + hipLaunchKernelStruct_t20 hipLaunchKernelStruct_h20 = + // out of order initalization + {.name = 'A', .rank = 2, .age = 42}; + bool *result_d20, *result_h20; + #if ENABLE_OUT_OF_ORDER_INITIALIZATION + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc20), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h20, result_d); + ResultValidation(); + #endif + + // Test: Passing the different hipLaunchParm options: + float* Ad; + hipMalloc((void**)&Ad, 1024); + hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), size_t(1024), 1, 0, 0, Ad); + hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), 1024, dim3(1), 0, 0, Ad); + hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), dim3(1024), 1, 0, 0, Ad); + hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), dim3(1024), dim3(1), 0, 0, Ad); + + // Test: Passing hipLaunchKernel inside another macro: float e0; - GPU_PRINT_TIME(hipLaunchKernel(vAdd, dim3(1024), + GPU_PRINT_TIME(hipLaunchKernelGGL(vAdd, dim3(1024), dim3(1), 0, 0, Ad), e0, j); - GPU_PRINT_TIME(WRAP(hipLaunchKernel(vAdd, dim3(1024), + GPU_PRINT_TIME(WRAP(hipLaunchKernelGGL(vAdd, dim3(1024), dim3(1), 0, 0, Ad)), e0, j); #ifdef EXTRA_PARENS_1 // Don't wrap hipLaunchKernel in extra set of parens: - GPU_PRINT_TIME((hipLaunchKernel(vAdd, dim3(1024), + GPU_PRINT_TIME((hipLaunchKernelGGL(vAdd, dim3(1024), dim3(1), 0, 0, Ad)), e0, j); #endif - MY_LAUNCH(hipLaunchKernel(vAdd, dim3(1024), dim3(1), + MY_LAUNCH(hipLaunchKernelGGL(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCall"); float* A; @@ -464,25 +779,12 @@ int main() { #ifdef EXTRA_PARENS_2 // MY_LAUNCH_WITH_PAREN wraps cmd in () which can cause issues. - MY_LAUNCH_WITH_PAREN(hipLaunchKernel(vAdd, dim3(1024), + MY_LAUNCH_WITH_PAREN(hipLaunchKernelGGL(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCall"); #endif - hipFree((void **)&result_h1); - hipFree((void **)&result_d1); - hipFree((void **)&result_h2); - hipFree((void **)&result_d2); - hipFree((void **)&result_h3); - hipFree((void **)&result_d3); - hipFree((void **)&result_h4); - hipFree((void **)&result_d4); - hipFree((void **)&result_h5); - hipFree((void **)&result_d5); - hipFree((void **)&result_h6); - hipFree((void **)&result_d6); - hipFree((void **)&result_h7); - hipFree((void **)&result_d7); - hipFree((void **)&result_h8); - hipFree((void **)&result_d8); + HIPCHECK(hipHostFree(result_h)); + HIPCHECK(hipFree(result_d)); + passed(); -} +} \ No newline at end of file From 4c707232b906e483f57c82a7df4c2cdaef5b631e Mon Sep 17 00:00:00 2001 From: Srinivasuluch Date: Wed, 18 Jul 2018 12:41:51 +0530 Subject: [PATCH 11/12] Update hipLaunchParm.cpp submitting changes as per Sam review/suggestions --- tests/src/kernel/hipLaunchParm.cpp | 162 +++++++++++++++++++++++++---- 1 file changed, 141 insertions(+), 21 deletions(-) diff --git a/tests/src/kernel/hipLaunchParm.cpp b/tests/src/kernel/hipLaunchParm.cpp index 420a74e7fa..0dd30c71c9 100644 --- a/tests/src/kernel/hipLaunchParm.cpp +++ b/tests/src/kernel/hipLaunchParm.cpp @@ -28,16 +28,25 @@ THE SOFTWARE. #include "test_common.h" // Memory alignment is broken -#define ENABLE_ALIGNMENT_TEST 0 -// Accessing struct class object from device is broken -#define ENABLE_CLASS_OBJ_ACCESS 0 +// Update: with latest changes the aligment is working fine, hence enabled +#define ENABLE_ALIGNMENT_TEST_SMALL_BAR 1 + +// Packed member atribute broken +#define ENABLE_PACKED_TEST 0 + +// Update: with latest changes struct class object +// from device is working fine, hence enabled +#define ENABLE_CLASS_OBJ_ACCESS 1 + // accessing dynamic/heap memory from device is broken #define ENABLE_HEAP_MEMORY_ACCESS 0 -// STL implementation broken, accessing dynamic/heap memory fail -// broken on hcc, working fine on hip-clang -#define ENABLE_USER_STL 0 -// out of order initalization is broken -#define ENABLE_OUT_OF_ORDER_INITIALIZATION 0 + +// Update: with latest changes it's working hence enabled +#define ENABLE_USER_STL 1 + +// Update: with latest changes it's working hence enabled +#define ENABLE_OUT_OF_ORDER_INITIALIZATION 1 + // Direct initialization of struct broken, // ip_d9 is a pointer, uint_t*, hipLaunchKernelStruct_h9 = {'c', ip_d9}; #define ENABLE_DECLARE_INITIALIZATION_POINTER 0 @@ -56,12 +65,20 @@ static hipError_t hipMemsetError = hipMemset(result_d, static void ResultValidation() { hipMemcpy(result_h, result_d, BLOCK_DIM_SIZE*sizeof(bool), hipMemcpyDeviceToHost); - for (int k = 0; k < BLOCK_DIM_SIZE; ++k) - HIPASSERT(result_h[k] == true); + for (int k = 0; k < BLOCK_DIM_SIZE; ++k) { + HIPASSERT(result_h[k] == true); + } + return; +} + +// Segregating the reset part as it was causing a problem when i put inside +// ResultValidation() function, the memory was not reset correctly for the +// tests which were disabled. +static void ResetValidationMem() { // reset the memory to false to reuse it. hipMemset(result_d, false, BLOCK_DIM_SIZE); - hipMemset(result_h, false, BLOCK_DIM_SIZE); +// hipMemset(result_h, false, BLOCK_DIM_SIZE); return; } @@ -125,6 +142,22 @@ typedef struct hipLaunchKernelStruct8 { bool b; }__attribute__((packed, aligned(4))) hipLaunchKernelStruct_t8; +// This test is to verify struct with packed, no alignment as Sam suggested +// size should be 4Bytes, right now it's broken on hcc & hip-clang +typedef struct hipLaunchKernelStruct8A { + char c1; + short int si; + bool b; +}__attribute__((packed)) hipLaunchKernelStruct_t8A; + +// This test is to verify struct with alignment, no packing as Sam suggested +// size should be 8Bytes as no packing, right now it's broken on hcc & hip-clang +typedef struct hipLaunchKernelStruct8B { + char c1; + short int si; + bool b; +}__attribute__((aligned(8))) hipLaunchKernelStruct_t8B; + // This test is to verify const struct object typedef struct hipLaunchKernelStruct9 { char c1; @@ -172,7 +205,7 @@ typedef struct hipLaunchKernelStruct14 { // the heap memory will be accessed from device typedef struct hipLaunchKernelStruct15 { char c1; - int* heapmem = new int[BLOCK_DIM_SIZE]; + int* heapmem; // allocated using hipMalloc() } hipLaunchKernelStruct_t15; // This test is to verify simple template struct @@ -344,7 +377,41 @@ __global__ void hipLaunchKernelStructFunc8( int *p = (int*)(&hipLaunchKernelStruct_); result_d8[x] = ((hipLaunchKernelStruct_.c1 == 'c') && (hipLaunchKernelStruct_.si == 1) - && ((size_t(p))%4 ==0) ); + && ((size_t(p))%4 ==0) + && (sizeof(hipLaunchKernelStruct_) == 4)); +} + +// Passing struct which is packed only, as Sam suggested, should be 4Bytes +// set the result_d8A to true if condition met +__global__ void hipLaunchKernelStructFunc8A( + hipLaunchKernelStruct_t8A hipLaunchKernelStruct_, + bool* result_d8A) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + // this is packed struct + // the address will not be aglined in this case hence condition removed + // only sizeof(hipLaunchKernelStruct_) will be valided + result_d8A[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (hipLaunchKernelStruct_.si == 1) + && (sizeof(hipLaunchKernelStruct_) == 4)); +} + +// Passing struct which is aligned(4) only, as Sam suggested +// , size should be 8Bytes, set the result_d8B to true if condition met +__global__ void hipLaunchKernelStructFunc8B( + hipLaunchKernelStruct_t8B hipLaunchKernelStruct_, + bool* result_d8B) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // set the result to true if the condition met + // get the address of the xth element, struct[x], + // size_t(p)%4 will be 0 if aligned to 4Byte address space + int *p = (int*)(&hipLaunchKernelStruct_); + result_d8B[x] = ((hipLaunchKernelStruct_.c1 == 'c') + && (hipLaunchKernelStruct_.si == 1) + && ((size_t(p))%8 == 0) + && (sizeof(hipLaunchKernelStruct_) == 8)); } // Passing struct with uint pointer object to a hipLaunchKernelGGL() @@ -474,7 +541,7 @@ __global__ void hipLaunchKernelStructFunc20( int x = blockIdx.x * blockDim.x + threadIdx.x; // accessing struct members in order - result_d20[x] = (hipLaunchKernelStruct_.name = 'A' + result_d20[x] = (hipLaunchKernelStruct_.name == 'A' // strcmp(hipLaunchKernelStruct_.name, "AMD") -> strcmp is not broken && hipLaunchKernelStruct_.age == 42 && hipLaunchKernelStruct_.rank == 2); @@ -520,6 +587,7 @@ int main() { HIPASSERT(hipMemsetError == hipSuccess); // Test: Passing Struct type, check access from device. + ResetValidationMem(); hipLaunchKernelStruct_t1 hipLaunchKernelStruct_h1; hipLaunchKernelStruct_h1.li = 1; hipLaunchKernelStruct_h1.lf = 1.0; @@ -531,6 +599,7 @@ int main() { ResultValidation(); // Test: Passing Struct type, checks padding + ResetValidationMem(); hipLaunchKernelStruct_t2 hipLaunchKernelStruct_h2; hipLaunchKernelStruct_h2.c1 = 'a'; hipLaunchKernelStruct_h2.l1 = 1.0; @@ -544,6 +613,7 @@ int main() { ResultValidation(); // Test: Passing Struct type, checks padding, assigning integer to a char + ResetValidationMem(); hipLaunchKernelStruct_t3 hipLaunchKernelStruct_h3; hipLaunchKernelStruct_h3.bf1 = 1; hipLaunchKernelStruct_h3.bf2 = 1; @@ -559,6 +629,7 @@ int main() { ResultValidation(); // Test: Passing empty struct + ResetValidationMem(); hipLaunchKernelStruct_t4 hipLaunchKernelStruct_h4; hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc4), dim3(BLOCK_DIM_SIZE), @@ -567,6 +638,7 @@ int main() { ResultValidation(); // Test: Passing struct with pointer object to a hipLaunchKernelGGL() + ResetValidationMem(); hipLaunchKernelStruct_t5 hipLaunchKernelStruct_h5; char* cp_d5; // This is passed as pointer to struct member // allocating memory for char pointer on device @@ -581,6 +653,7 @@ int main() { ResultValidation(); // Test: Passing struct with aligned(8) + ResetValidationMem(); hipLaunchKernelStruct_t6 hipLaunchKernelStruct_h6; hipLaunchKernelStruct_h6.c1 = 'c'; hipLaunchKernelStruct_h6.si = 1; @@ -589,16 +662,17 @@ int main() { dim3(1), 0, 0, hipLaunchKernelStruct_h6, result_d); // alignment is broken hence disabled the validation part - #if ENABLE_ALIGNMENT_TEST + #if ENABLE_ALIGNMENT_TEST_SMALL_BAR ResultValidation(); #endif // Test: Passing struct with aligned(16) + ResetValidationMem(); hipLaunchKernelStruct_t7 hipLaunchKernelStruct_h7; hipLaunchKernelStruct_h7.c1 = 'c'; hipLaunchKernelStruct_h7.si = 1; - #if ENABLE_ALIGNMENT_TEST // This is broken on small bar + #if ENABLE_ALIGNMENT_TEST_SMALL_BAR // This is broken on small bar hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc7), dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h7, @@ -606,7 +680,8 @@ int main() { ResultValidation(); #endif - // Test: Passing struct with packed aligned to 6Bytes + // Test: Passing struct with packed aligned to 4Bytes + ResetValidationMem(); hipLaunchKernelStruct_t8 hipLaunchKernelStruct_h8; hipLaunchKernelStruct_h8.c1 = 'c'; hipLaunchKernelStruct_h8.si = 1; @@ -614,13 +689,41 @@ int main() { dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h8, result_d); - // alignment is broken hence disabled the validation part - #if ENABLE_ALIGNMENT_TEST + // packed member broken on large and small bar setup. + #if ENABLE_PACKED_TEST ResultValidation(); #endif + // Test: Passing struct with packed to 4Bytes + ResetValidationMem(); + hipLaunchKernelStruct_t8A hipLaunchKernelStruct_h8A; + hipLaunchKernelStruct_h8A.c1 = 'c'; + hipLaunchKernelStruct_h8A.si = 1; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc8A), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h8A, + result_d); + // packed member broken on large and small bar setup. + #if ENABLE_PACKED_TEST + ResultValidation(); + #endif + + // Test: Passing struct with aligned(4) to 4Bytes, size is 8Bytes + ResetValidationMem(); + hipLaunchKernelStruct_t8B hipLaunchKernelStruct_h8B; + hipLaunchKernelStruct_h8B.c1 = 'c'; + hipLaunchKernelStruct_h8B.si = 1; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc8B), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h8B, + result_d); + // alignment is broken hence disabled the validation part + #if ENABLE_ALIGNMENT_TEST_SMALL_BAR + ResultValidation(); + #endif // Test: Passing const struct object to a hipLaunchKernelGGL() + ResetValidationMem(); uint32_t* ip_d9; // allocating memory for char pointer on device HIPCHECK(hipMalloc((void**)&ip_d9, sizeof(uint32_t))); @@ -637,6 +740,7 @@ int main() { // Test: Passing struct with uintN_t as member variables + ResetValidationMem(); hipLaunchKernelStruct_t10 hipLaunchKernelStruct_h10; hipLaunchKernelStruct_h10.u64 = UINT64_MAX; hipLaunchKernelStruct_h10.u32 = 1; @@ -649,6 +753,7 @@ int main() { // Test: Passing struct with uintN_t as member variables + ResetValidationMem(); hipLaunchKernelStruct_t11 hipLaunchKernelStruct_h11; hipLaunchKernelStruct_h11.i1 = 1; hipLaunchKernelStruct_h11.vint = 0; @@ -659,6 +764,7 @@ int main() { ResultValidation(); // Test: Passing struct with simple class object + ResetValidationMem(); hipLaunchKernelStruct_t12 hipLaunchKernelStruct_h12; hipLaunchKernelStruct_h12.c1 = 'c'; hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc12), @@ -671,6 +777,7 @@ int main() { #endif // Test: Passing struct with simple __device__ func() + ResetValidationMem(); hipLaunchKernelStruct_t13 hipLaunchKernelStruct_h13; hipLaunchKernelStruct_h13.i1 = 1; hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc13), @@ -680,6 +787,7 @@ int main() { ResultValidation(); // Test: Passing struct with array variable, write to from device + ResetValidationMem(); hipLaunchKernelStruct_t14 hipLaunchKernelStruct_h14; hipLaunchKernelStruct_h14.readint = 1; hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc14), @@ -690,9 +798,16 @@ int main() { // Test: Passing struct with heap memory, read to from device + ResetValidationMem(); hipLaunchKernelStruct_t15 hipLaunchKernelStruct_h15; hipLaunchKernelStruct_h15.c1 = 'c'; - #if ENABLE_HEAP_MEMORY_ACCESS // causing page fault here,validation failed + + #if ENABLE_HEAP_MEMORY_ACCESS // causing page fault here, + // on small bar set + HIPCHECK(hipMalloc(&hipLaunchKernelStruct_h15.heapmem, + BLOCK_DIM_SIZE*sizeof(int))); + HIPCHECK(hipMemset(&hipLaunchKernelStruct_h15.heapmem, + 0, BLOCK_DIM_SIZE)); hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc15), dim3(BLOCK_DIM_SIZE), dim3(1), 0, 0, hipLaunchKernelStruct_h15, @@ -701,6 +816,7 @@ int main() { #endif // Test: Passing simple template struct + ResetValidationMem(); hipLaunchKernelStruct_t16 hipLaunchKernelStruct_h16; hipLaunchKernelStruct_h16.t1 = 'c'; hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc16), @@ -710,6 +826,7 @@ int main() { ResultValidation(); // Test: Passing simple explicit template struct + ResetValidationMem(); hipLaunchKernelStruct_t17 hipLaunchKernelStruct_h17; hipLaunchKernelStruct_h17.t1 = 1; hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc17), @@ -719,6 +836,7 @@ int main() { ResultValidation(); // Test: Passing struct with simple __device__ func() to struct memory + ResetValidationMem(); hipLaunchKernelStruct_t18 hipLaunchKernelStruct_h18; hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc18), dim3(BLOCK_DIM_SIZE), @@ -727,6 +845,7 @@ int main() { ResultValidation(); // Test: Passing user defined stack, + ResetValidationMem(); hipLaunchKernelStruct_t19 hipLaunchKernelStruct_h19; hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc19), dim3(BLOCK_DIM_SIZE), @@ -738,6 +857,7 @@ int main() { // Test: Passing struct which is initiazed out of order // accessing same elements in order from device + ResetValidationMem(); hipLaunchKernelStruct_t20 hipLaunchKernelStruct_h20 = // out of order initalization {.name = 'A', .rank = 2, .age = 42}; @@ -787,4 +907,4 @@ int main() { HIPCHECK(hipFree(result_d)); passed(); -} \ No newline at end of file +} From b37c56efedb6899ab1ab4697a8f15246a6c1b436 Mon Sep 17 00:00:00 2001 From: Srinivasuluch Date: Fri, 20 Jul 2018 12:38:06 +0530 Subject: [PATCH 12/12] Update hipLaunchParm.cpp Added struct bit field test, test number#21 --- tests/src/kernel/hipLaunchParm.cpp | 44 ++++++++++++++++++++++++++++-- 1 file changed, 41 insertions(+), 3 deletions(-) diff --git a/tests/src/kernel/hipLaunchParm.cpp b/tests/src/kernel/hipLaunchParm.cpp index 0dd30c71c9..befdcc656d 100644 --- a/tests/src/kernel/hipLaunchParm.cpp +++ b/tests/src/kernel/hipLaunchParm.cpp @@ -51,6 +51,9 @@ THE SOFTWARE. // ip_d9 is a pointer, uint_t*, hipLaunchKernelStruct_h9 = {'c', ip_d9}; #define ENABLE_DECLARE_INITIALIZATION_POINTER 0 +// Bit fields are broken +#define ENABLE_BIT_FIELDS 0 + static const int BLOCK_DIM_SIZE = 1024; // allocate memory on device and host for result validation @@ -78,7 +81,7 @@ static void ResultValidation() { static void ResetValidationMem() { // reset the memory to false to reuse it. hipMemset(result_d, false, BLOCK_DIM_SIZE); -// hipMemset(result_h, false, BLOCK_DIM_SIZE); + hipMemset(result_h, false, BLOCK_DIM_SIZE); return; } @@ -237,8 +240,10 @@ typedef struct hipLaunchKernelStruct19 { stackNode_t* stack = NULL; unsigned int size_ = 0; void pushMe(int value) { // not a device function, setting from host - stackNode_t* newNode = new stackNode_t; - newNode->data = value; + stackNode_t* newNode; + hipMalloc((void**)&newNode, sizeof(stackNode_t)); + hipMemset(&newNode->data, value, sizeof(stackNode_t)); + //newNode->data = value; ++size_; if (stack == NULL) { stack = newNode; @@ -269,6 +274,13 @@ typedef struct hipLaunchKernelStruct20 { int rank; } hipLaunchKernelStruct_t20; +// This test is to verify bit fields operations +// the size should be 1Bytes +typedef struct hipLaunchKernelStruct21 { + int i : 3; // limiting bits to 3 + int j : 2; // limiting bits to 2 +} hipLaunchKernelStruct_t21; + // Passing struct to a hipLaunchKernelGGL(), // read and write into the same struct __global__ void hipLaunchKernelStructFunc1( @@ -547,6 +559,18 @@ __global__ void hipLaunchKernelStructFunc20( && hipLaunchKernelStruct_.rank == 2); } +// Passing struct with bit fields +__global__ void hipLaunchKernelStructFunc21( + hipLaunchKernelStruct_t21 hipLaunchKernelStruct_, + bool* result_d21) { + int x = blockIdx.x * blockDim.x + threadIdx.x; + + // accessing struct members in order + result_d21[x] = (hipLaunchKernelStruct_.i == 2 + && hipLaunchKernelStruct_.j == 0 + && (sizeof(hipLaunchKernelStruct_) == 1)); +} + __global__ void vAdd(float* a) {} //--- @@ -869,6 +893,20 @@ int main() { ResultValidation(); #endif + // Test: Passing struct with bit fields operation + // accessing same elements in order from device + ResetValidationMem(); + hipLaunchKernelStruct_t21 hipLaunchKernelStruct_h21 = + // out of order initalization + {2,0}; + bool *result_d21, *result_h21; + hipLaunchKernelGGL(HIP_KERNEL_NAME(hipLaunchKernelStructFunc21), + dim3(BLOCK_DIM_SIZE), + dim3(1), 0, 0, hipLaunchKernelStruct_h21, result_d); + #if ENABLE_BIT_FIELDS + ResultValidation(); + #endif + // Test: Passing the different hipLaunchParm options: float* Ad; hipMalloc((void**)&Ad, 1024);