From 299d2358dd8f2c0f1259f3cef449bd2b778fdde0 Mon Sep 17 00:00:00 2001 From: Jin Z <5zj@equinox.ftpn.ornl.gov> Date: Tue, 18 Apr 2023 11:56:10 -0400 Subject: [PATCH] [aligned-types-] replace uint4_aligned_2 with uint8_aligned; add the tests of uint8 aligned/misaligned types --- aligned-types-cuda/main.cu | 14 +++++++++++--- aligned-types-hip/main.cu | 16 ++++++++++++---- aligned-types-omp/main.cpp | 13 ++++++++++--- 3 files changed, 33 insertions(+), 10 deletions(-) diff --git a/aligned-types-cuda/main.cu b/aligned-types-cuda/main.cu index fdb398a7c..2dc434f78 100644 --- a/aligned-types-cuda/main.cu +++ b/aligned-types-cuda/main.cu @@ -52,6 +52,11 @@ typedef struct unsigned int r, g, b, a; } uint4_misaligned; +typedef struct +{ + uint4_misaligned c1, c2; +} +uint8_misaligned; //////////////////////////////////////////////////////////////////////////////// @@ -97,7 +102,7 @@ typedef struct __align__(16) { uint4_aligned c1, c2; } -uint4_aligned_2; +uint8_aligned; @@ -303,8 +308,11 @@ int main(int argc, char **argv) printf("uint4_aligned...\n"); nTotalFailures += runTest(d_idata, d_odata, 16, MemorySize); - printf("uint4_aligned_2...\n"); - nTotalFailures += runTest(d_idata, d_odata, 32, MemorySize); + printf("uint8_misaligned...\n"); + nTotalFailures += runTest(d_idata, d_odata, 32, MemorySize); + + printf("uint8_aligned...\n"); + nTotalFailures += runTest(d_idata, d_odata, 32, MemorySize); printf("\n[alignedTypes] -> Test Results: %d Failures\n", nTotalFailures); diff --git a/aligned-types-hip/main.cu b/aligned-types-hip/main.cu index 50f054a5f..dfa8d7212 100644 --- a/aligned-types-hip/main.cu +++ b/aligned-types-hip/main.cu @@ -52,6 +52,11 @@ typedef struct unsigned int r, g, b, a; } uint4_misaligned; +typedef struct +{ + uint4_misaligned c1, c2; +} +uint8_misaligned; //////////////////////////////////////////////////////////////////////////////// @@ -97,7 +102,7 @@ typedef struct __align__(16) { uint4_aligned c1, c2; } -uint4_aligned_2; +uint8_aligned; @@ -211,7 +216,7 @@ template int runTest( dim3 block (256); for (int i = 0; i < NUM_ITERATIONS; i++) { - hipLaunchKernelGGL(HIP_KERNEL_NAME(testKernel), grid, block, 0, 0, + testKernel<<>>( (TData *)d_odata, (TData *)d_idata, numElements @@ -303,8 +308,11 @@ int main(int argc, char **argv) printf("uint4_aligned...\n"); nTotalFailures += runTest(d_idata, d_odata, 16, MemorySize); - printf("uint4_aligned_2...\n"); - nTotalFailures += runTest(d_idata, d_odata, 32, MemorySize); + printf("uint8_misaligned...\n"); + nTotalFailures += runTest(d_idata, d_odata, 32, MemorySize); + + printf("uint8_aligned...\n"); + nTotalFailures += runTest(d_idata, d_odata, 32, MemorySize); printf("\n[alignedTypes] -> Test Results: %d Failures\n", nTotalFailures); diff --git a/aligned-types-omp/main.cpp b/aligned-types-omp/main.cpp index 211f9af61..474846044 100644 --- a/aligned-types-omp/main.cpp +++ b/aligned-types-omp/main.cpp @@ -52,6 +52,10 @@ typedef struct unsigned int r, g, b, a; } uint4_misaligned; +typedef struct +{ + uint4_misaligned c1, c2; +} uint8_misaligned; //////////////////////////////////////////////////////////////////////////////// @@ -97,7 +101,7 @@ typedef struct __attribute__((__aligned__(16))) { uint4_aligned c1, c2; } -uint4_aligned_2; +uint8_aligned; @@ -280,8 +284,11 @@ int main(int argc, char **argv) printf("uint4_aligned...\n"); nTotalFailures += runTest(d_idata, d_odata, 16, MemorySize); - printf("uint4_aligned_2...\n"); - nTotalFailures += runTest(d_idata, d_odata, 32, MemorySize); + printf("uint8_misaligned...\n"); + nTotalFailures += runTest(d_idata, d_odata, 32, MemorySize); + + printf("uint8_aligned...\n"); + nTotalFailures += runTest(d_idata, d_odata, 32, MemorySize); printf("\n[alignedTypes] -> Test Results: %d Failures\n", nTotalFailures);