memset_offset_current.cpp

#include <hip/hip_runtime.h>
#include <assert.h>
#include <cstring>
#include <iostream>
#include <set>
#include <map>
#include <cstdlib>
#include <string>

#define HIP_CHECK(error)                                                                           \
  {                                                                                                \
    hipError_t localError = error;                                                                 \
    if ((localError != hipSuccess)) {                                                              \
      printf("HIP error %d at %s:%d\n", localError, __FILE__, __LINE__);                           \
      assert(false);                                                                               \
    }                                                                                              \
  }

// Poison value: chosen to be distinct from any memsetval we use
#define POISON_VALUE 0xDE

__global__ void memset_(char* dst, char value, size_t N, unsigned* wid, unsigned* xcc) {
    size_t i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        unsigned smid = __smid();
        wid[i] = smid;
        xcc[i] = (smid >> 9) & 0xF;
        dst[i] = value;
    }
}

__global__ void verify_value(char* dst, char value, size_t N, unsigned* wid, unsigned* xcc, bool* mismatch) {
    size_t i = blockIdx.x * blockDim.x + threadIdx.x;
    size_t idx = N - i - 1;
    if (idx < N && idx >= 0) {
        unsigned smid = __smid();
        wid[i] = smid;
        xcc[i] = (smid >> 9) & 0xF;
        dst[i] = value;
        if (dst[idx] != value) {
            //printf("Mismatch at i=%zu: expected 0x%02X got 0x%02X\n",
            //       i, value, dst[i]);
            mismatch[idx] = true;
            //abort();
        }
    }
}

__global__ void memset_xcd(char* dst, char value, size_t N, size_t xcd, unsigned* xcd_count) {
#if 0
    unsigned smid = __smid();
    auto xcc = (smid >> 9) & 0xF;  
    if (xcc == xcd) {
        __shared__ size_t c;
        if (threadIdx.x == 0) {
            c = atomicAdd(xcd_count, 1);
        }
        __syncthreads();
        if (c == 0 && threadIdx.x == 0) {
#if 1
            for (size_t i = 0; i < N; i ++) {
                dst[i] = value;
            }
            #endif
        }
    }
#endif
}

static bool testhipMemset(char memsetval, size_t numElements, size_t offset = 0) {
  if (numElements <= offset) {
    printf("Error: numElements (%zu) must be greater than offset (%zu)\n", numElements, offset);
    return false;
  }
  size_t Nbytes = numElements * sizeof(char);
  bool testResult = true;

  printf("\n=== Test: memsetval=0x%02X numElements=%zu offset=%zu ===\n",
         (unsigned char)memsetval, numElements, offset);

  char* A_d = nullptr;
  HIP_CHECK(hipMalloc(&A_d, Nbytes));
  char* A_h = reinterpret_cast<char*>(malloc(Nbytes));
  assert(A_h != nullptr);

  // ---------------------------------------------------------------
  // 1) Poison device memory with a known value from CPU
  //    Fill host buffer with POISON_VALUE, then copy to device.
  //    This way, if the kernel fails to write a byte, we'll see
  //    POISON_VALUE instead of uninitialized garbage.
  // ---------------------------------------------------------------
  ::memset(A_h, POISON_VALUE, Nbytes);
  HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
  printf("Poisoned device memory with 0x%02X (%zu bytes)\n", POISON_VALUE, Nbytes);

  // ---------------------------------------------------------------
  // 2) Verify poison landed correctly by reading back immediately
  // ---------------------------------------------------------------
  char* verify_h = reinterpret_cast<char*>(malloc(Nbytes));
  assert(verify_h != nullptr);
  HIP_CHECK(hipMemcpy(verify_h, A_d, Nbytes, hipMemcpyDeviceToHost));
  bool poisonOk = true;
  for (size_t i = 0; i < numElements; i++) {
    if (verify_h[i] != (char)POISON_VALUE) {
      printf("Poison verify FAILED at i=%zu: expected 0x%02X got 0x%02X\n",
             i, POISON_VALUE,
             (unsigned)(unsigned char)verify_h[i]);
      poisonOk = false;
      break;
    }
  }
  if (poisonOk) {
    printf("Poison verify OK: device buffer confirmed 0x%02X before kernel\n", POISON_VALUE);
  }
  free(verify_h);

  // ---------------------------------------------------------------
  // 3) Launch custom memset kernel with XCC tracking
  // ---------------------------------------------------------------
  size_t Nmemset = numElements - offset;
  const unsigned blockSize = 256;
  int maxGridX = 0;
  HIP_CHECK(hipDeviceGetAttribute(&maxGridX, hipDeviceAttributeMaxGridDimX, 0));
  unsigned gridSize = (unsigned)((Nmemset + blockSize - 1) / blockSize);
  if (gridSize == 0) gridSize = 1;
  if (gridSize > (unsigned)maxGridX) gridSize = (unsigned)maxGridX;

  unsigned* wid_d = nullptr;
  unsigned* xcc_d = nullptr;
  HIP_CHECK(hipMalloc(&wid_d, Nmemset * sizeof(unsigned)));
  HIP_CHECK(hipMalloc(&xcc_d, Nmemset * sizeof(unsigned)));

  printf("Launching kernel: gridSize=%u blockSize=%u Nmemset=%zu\n", gridSize, blockSize, Nmemset);
  memset_<<<gridSize, blockSize>>>(A_d + offset, memsetval, Nmemset, wid_d, xcc_d);
  bool* mismatch_d = nullptr;
  HIP_CHECK(hipMalloc(&mismatch_d, Nmemset * sizeof(bool)));
  HIP_CHECK(hipMemset(mismatch_d, 0, Nmemset * sizeof(bool)));
  HIP_CHECK(hipGetLastError());
  HIP_CHECK(hipDeviceSynchronize());

#if 0  /**/
  printf("Checking values inside A_d directly on CPU\n\n\n\n\n");
  for (size_t i = 0; i < Nmemset; i++) {
    if (A_d[i + offset] != memsetval) {
      printf("Mismatch at i=%zu: expected 0x%02X got 0x%02X\n",
             i, memsetval, A_d[i + offset]);\
    }
  }
#endif


#if 0
  printf("Checking values inside A_d directly on GPU\n\n\n\n\n");
  verify_value<<<gridSize, blockSize>>>(A_d + offset, memsetval, Nmemset, wid_d, xcc_d, mismatch_d);
#endif

  HIP_CHECK(hipGetLastError());
  HIP_CHECK(hipDeviceSynchronize());

  for (size_t i = 0; i < Nmemset; i++) {
    if (mismatch_d[i]) {
      printf("Mismatch at i=%zu: expected 0x%02X got 0x%02X\n",
             i, memsetval, A_d[i + offset]);
      }
  }
  HIP_CHECK(hipFree(mismatch_d));
  // ---------------------------------------------------------------
  // 4) Copy back XCC tracking data
  // ---------------------------------------------------------------
  unsigned* wid_h = reinterpret_cast<unsigned*>(malloc(Nmemset * sizeof(unsigned)));
  unsigned* xcc_h = reinterpret_cast<unsigned*>(malloc(Nmemset * sizeof(unsigned)));
  assert(wid_h != nullptr && xcc_h != nullptr);
  HIP_CHECK(hipMemcpy(wid_h, wid_d, Nmemset * sizeof(unsigned), hipMemcpyDeviceToHost));
  HIP_CHECK(hipMemcpy(xcc_h, xcc_d, Nmemset * sizeof(unsigned), hipMemcpyDeviceToHost));

  // ---------------------------------------------------------------
  // 5) Copy back device buffer and check from CPU
  // ---------------------------------------------------------------
  ::memset(A_h, 0, Nbytes);  // clear host buffer before readback
  HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));

  // Check region before offset still has poison (kernel should not have touched it)
  for (size_t i = 0; i < offset; i++) {
    if (A_h[i] != (char)POISON_VALUE) {
      printf("WARNING: pre-offset region corrupted at i=%zu: expected poison 0x%02X got 0x%02X\n",
             i, POISON_VALUE, (unsigned)(unsigned char)A_h[i]);
    }
  }

  // ---------------------------------------------------------------
  // 6) Check kernel-written region and collect XCC stats
  // ---------------------------------------------------------------
  std::set<unsigned> xcc_set;
  std::map<unsigned, size_t> xcc_counts;       // xcc -> number of elements written
  std::map<unsigned, size_t> xcc_fail_counts;  // xcc -> number of mismatches
  size_t mismatchCount = 0;
  size_t mismatchPoisonCount = 0;  // mismatches that are still the poison value

  printf("Checking values inside after Memcpy\n\n\n\n\n");
  for (size_t i = offset; i < numElements; i++) {
    size_t ki = i - offset;  // index into kernel arrays
    unsigned this_xcc = xcc_h[ki];
    xcc_set.insert(this_xcc);
    xcc_counts[this_xcc]++;

    if (A_h[i] != memsetval) {
      testResult = false;
      mismatchCount++;
      xcc_fail_counts[this_xcc]++;

      if (A_h[i] == (char)POISON_VALUE) {
        mismatchPoisonCount++;
      }

      if (mismatchCount <= 20) {  // print first 20 mismatches
        std::cout << "Mismatch data : "
                  << "numElements: " << numElements << " "
                  << "offset: " << offset << " "
                  << "i: " << i << " "
                  << "A_h[i]: 0x" << std::hex << std::uppercase
                  << static_cast<unsigned>(static_cast<unsigned char>(A_h[i]))
                  << std::dec << std::nouppercase << " "
                  << "memsetval: 0x" << std::hex << std::uppercase
                  << static_cast<unsigned>(static_cast<unsigned char>(memsetval))
                  << std::dec << std::nouppercase << " "
                  << "smid: " << wid_h[ki] << " "
                  << "xcc: " << this_xcc
                  << (A_h[i] == (char)POISON_VALUE ? " [STILL POISON]" : " [CORRUPTED]")
                  << std::endl;
      }
    }
  }

  // ---------------------------------------------------------------
  // 7) Print XCC summary (always, even in passing case)
  // ---------------------------------------------------------------
  printf("\n--- XCC Summary ---\n");
  printf("XCCs seen (count=%zu): {", xcc_set.size());
  for (auto it = xcc_set.begin(); it != xcc_set.end(); ++it) {
    if (it != xcc_set.begin()) printf(", ");
    printf("%u", *it);
  }
  printf("}\n");

  // Check XCC mask: XCD 0-3 means XCC values 0-3 (mask 0xF)
  bool allOnOneAid = true;
  for (unsigned xcc_val : xcc_set) {
    if (xcc_val > 3) {
      allOnOneAid = false;
      break;
    }
  }
  printf("All XCCs in XCD 0-3 (mask 0xF, one AID): %s\n", allOnOneAid ? "YES" : "NO");

  printf("\nPer-XCC element counts:\n");
  for (auto& [xcc_val, count] : xcc_counts) {
    printf("  XCC %2u: %zu elements", xcc_val, count);
    if (xcc_fail_counts.count(xcc_val)) {
      printf("  *** %zu MISMATCHES ***", xcc_fail_counts[xcc_val]);
    }
    printf("\n");
  }

  if (mismatchCount > 0) {
    printf("\nTotal mismatches: %zu / %zu (%.4f%%)\n",
           mismatchCount, Nmemset, 100.0 * mismatchCount / Nmemset);
    printf("  Of which still poison (0x%02X): %zu\n", POISON_VALUE, mismatchPoisonCount);
    printf("  Of which corrupted (other value): %zu\n", mismatchCount - mismatchPoisonCount);
  } else {
    printf("\nResult: PASS (all %zu elements match 0x%02X)\n",
           Nmemset, (unsigned char)memsetval);
  }
  printf("-------------------\n");

  HIP_CHECK(hipFree(A_d));
  HIP_CHECK(hipFree(wid_d));
  HIP_CHECK(hipFree(xcc_d));
  free(A_h);
  free(wid_h);
  free(xcc_h);
  return testResult;
}

static bool testhipMemsetXcd(char memsetval, size_t numElements, size_t xcd, size_t offset = 0) {
    if (numElements <= offset) {
      printf("Error: numElements (%zu) must be greater than offset (%zu)\n", numElements, offset);
      return false;
    }
    size_t Nbytes = numElements * sizeof(char);
    bool testResult = true;
  
    printf("\n=== Test: memsetval=0x%02X numElements=%zu offset=%zu xcd=%zu ===\n",
           (unsigned char)memsetval, numElements, offset, xcd);
  
    char* A_d = nullptr;
    HIP_CHECK(hipMalloc(&A_d, Nbytes));
    char* A_h = reinterpret_cast<char*>(malloc(Nbytes));
    assert(A_h != nullptr);
  
    // ---------------------------------------------------------------
    // 1) Poison device memory with a known value from CPU
    //    Fill host buffer with POISON_VALUE, then copy to device.
    //    This way, if the kernel fails to write a byte, we'll see
    //    POISON_VALUE instead of uninitialized garbage.
    // ---------------------------------------------------------------
    ::memset(A_h, POISON_VALUE, Nbytes);
    HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
    printf("Poisoned device memory with 0x%02X (%zu bytes)\n", POISON_VALUE, Nbytes);
  
    // ---------------------------------------------------------------
    // 2) Verify poison landed correctly by reading back immediately
    // ---------------------------------------------------------------
    char* verify_h = reinterpret_cast<char*>(malloc(Nbytes));
    assert(verify_h != nullptr);
    HIP_CHECK(hipMemcpy(verify_h, A_d, Nbytes, hipMemcpyDeviceToHost));
    bool poisonOk = true;
    for (size_t i = 0; i < numElements; i++) {
      if (verify_h[i] != (char)POISON_VALUE) {
        printf("Poison verify FAILED at i=%zu: expected 0x%02X got 0x%02X\n",
               i, POISON_VALUE,
               (unsigned)(unsigned char)verify_h[i]);
        poisonOk = false;
        break;
      }
    }
    if (poisonOk) {
      printf("Poison verify OK: device buffer confirmed 0x%02X before kernel\n", POISON_VALUE);
    }
    free(verify_h);
  
    // ---------------------------------------------------------------
    // 3) Launch custom memset kernel with XCC tracking
    // ---------------------------------------------------------------
    size_t Nmemset = numElements - offset;
    const unsigned blockSize = 256;
    int maxGridX = 0;
    HIP_CHECK(hipDeviceGetAttribute(&maxGridX, hipDeviceAttributeMaxGridDimX, 0));
    unsigned gridSize = (unsigned)((Nmemset + blockSize - 1) / blockSize);
    if (gridSize == 0) gridSize = 1;
    if (gridSize > (unsigned)maxGridX) gridSize = (unsigned)maxGridX;
  
    unsigned* wid_d = nullptr;
    unsigned* xcc_d = nullptr;
    HIP_CHECK(hipMalloc(&wid_d, Nmemset * sizeof(unsigned)));
    HIP_CHECK(hipMalloc(&xcc_d, Nmemset * sizeof(unsigned)));

    unsigned* xcd_count_d = nullptr;
    HIP_CHECK(hipMalloc(&xcd_count_d, sizeof(unsigned)));
    HIP_CHECK(hipMemset(xcd_count_d, 0, sizeof(unsigned)));
  
    printf("Launching kernel: gridSize=%u blockSize=%u Nmemset=%zu xcd=%zu\n", gridSize, blockSize, Nmemset, xcd);
    memset_xcd<<<gridSize, blockSize>>>(A_d + offset, memsetval, Nmemset, xcd, xcd_count_d);

    HIP_CHECK(hipGetLastError());
    HIP_CHECK(hipDeviceSynchronize());

    bool* mismatch_d = nullptr;
    HIP_CHECK(hipMalloc(&mismatch_d, Nmemset * sizeof(bool)));
    HIP_CHECK(hipMemset(mismatch_d, 0, Nmemset * sizeof(bool)));
    HIP_CHECK(hipGetLastError());
    HIP_CHECK(hipDeviceSynchronize());
  
  #if 0
    printf("Checking values inside A_d directly on CPU\n\n\n\n\n");
    for (size_t i = 0; i < Nmemset; i++) {
      if (A_d[i + offset] != memsetval) {
        printf("Mismatch at i=%zu: expected 0x%02X got 0x%02X\n",
               i, memsetval, A_d[i + offset]);\
      }
    }
  #endif
  
    printf("Checking values inside A_d directly on GPU\n\n\n\n\n");
    verify_value<<<gridSize, blockSize>>>(A_d + offset, memsetval, Nmemset, wid_d, xcc_d, mismatch_d);
    HIP_CHECK(hipGetLastError());
    HIP_CHECK(hipDeviceSynchronize());
  
    for (size_t i = 0; i < Nmemset; i++) {
      if (mismatch_d[i]) {
        printf("Mismatch at i=%zu: expected 0x%02X got 0x%02X\n",
               i, memsetval, A_d[i + offset]);
        }
    }
    HIP_CHECK(hipFree(mismatch_d));
    // ---------------------------------------------------------------
    // 4) Copy back XCC tracking data
    // ---------------------------------------------------------------
    unsigned* wid_h = reinterpret_cast<unsigned*>(malloc(Nmemset * sizeof(unsigned)));
    unsigned* xcc_h = reinterpret_cast<unsigned*>(malloc(Nmemset * sizeof(unsigned)));
    assert(wid_h != nullptr && xcc_h != nullptr);
    HIP_CHECK(hipMemcpy(wid_h, wid_d, Nmemset * sizeof(unsigned), hipMemcpyDeviceToHost));
    HIP_CHECK(hipMemcpy(xcc_h, xcc_d, Nmemset * sizeof(unsigned), hipMemcpyDeviceToHost));
  
    // ---------------------------------------------------------------
    // 5) Copy back device buffer and check from CPU
    // ---------------------------------------------------------------
    ::memset(A_h, 0, Nbytes);  // clear host buffer before readback
    HIP_CHECK(hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost));
  
    // Check region before offset still has poison (kernel should not have touched it)
    for (size_t i = 0; i < offset; i++) {
      if (A_h[i] != (char)POISON_VALUE) {
        printf("WARNING: pre-offset region corrupted at i=%zu: expected poison 0x%02X got 0x%02X\n",
               i, POISON_VALUE, (unsigned)(unsigned char)A_h[i]);
      }
    }
  
    // ---------------------------------------------------------------
    // 6) Check kernel-written region and collect XCC stats
    // ---------------------------------------------------------------
    std::set<unsigned> xcc_set;
    std::map<unsigned, size_t> xcc_counts;       // xcc -> number of elements written
    std::map<unsigned, size_t> xcc_fail_counts;  // xcc -> number of mismatches
    size_t mismatchCount = 0;
    size_t mismatchPoisonCount = 0;  // mismatches that are still the poison value
  
    printf("Checking values inside after Memcpy\n\n\n\n\n");
    for (size_t i = offset; i < numElements; i++) {
      size_t ki = i - offset;  // index into kernel arrays
      unsigned this_xcc = xcc_h[ki];
      xcc_set.insert(this_xcc);
      xcc_counts[this_xcc]++;
  
      if (A_h[i] != memsetval) {
        testResult = false;
        mismatchCount++;
        xcc_fail_counts[this_xcc]++;
  
        if (A_h[i] == (char)POISON_VALUE) {
          mismatchPoisonCount++;
        }
  
        if (mismatchCount <= 20) {  // print first 20 mismatches
          std::cout << "Mismatch data : "
                    << "numElements: " << numElements << " "
                    << "offset: " << offset << " "
                    << "i: " << i << " "
                    << "A_h[i]: 0x" << std::hex << std::uppercase
                    << static_cast<unsigned>(static_cast<unsigned char>(A_h[i]))
                    << std::dec << std::nouppercase << " "
                    << "memsetval: 0x" << std::hex << std::uppercase
                    << static_cast<unsigned>(static_cast<unsigned char>(memsetval))
                    << std::dec << std::nouppercase << " "
                    << "smid: " << wid_h[ki] << " "
                    << "xcc: " << this_xcc
                    << (A_h[i] == (char)POISON_VALUE ? " [STILL POISON]" : " [CORRUPTED]")
                    << std::endl;
        }
      }
    }
  
    // ---------------------------------------------------------------
    // 7) Print XCC summary (always, even in passing case)
    // ---------------------------------------------------------------
    printf("\n--- XCC Summary ---\n");
    printf("XCCs seen (count=%zu): {", xcc_set.size());
    for (auto it = xcc_set.begin(); it != xcc_set.end(); ++it) {
      if (it != xcc_set.begin()) printf(", ");
      printf("%u", *it);
    }
    printf("}\n");
  
    // Check XCC mask: XCD 0-3 means XCC values 0-3 (mask 0xF)
    bool allOnOneAid = true;
    for (unsigned xcc_val : xcc_set) {
      if (xcc_val > 3) {
        allOnOneAid = false;
        break;
      }
    }
    printf("All XCCs in XCD 0-3 (mask 0xF, one AID): %s\n", allOnOneAid ? "YES" : "NO");
  
    printf("\nPer-XCC element counts:\n");
    for (auto& [xcc_val, count] : xcc_counts) {
      printf("  XCC %2u: %zu elements", xcc_val, count);
      if (xcc_fail_counts.count(xcc_val)) {
        printf("  *** %zu MISMATCHES ***", xcc_fail_counts[xcc_val]);
      }
      printf("\n");
    }
  
    if (mismatchCount > 0) {
      printf("\nTotal mismatches: %zu / %zu (%.4f%%)\n",
             mismatchCount, Nmemset, 100.0 * mismatchCount / Nmemset);
      printf("  Of which still poison (0x%02X): %zu\n", POISON_VALUE, mismatchPoisonCount);
      printf("  Of which corrupted (other value): %zu\n", mismatchCount - mismatchPoisonCount);
    } else {
      printf("\nResult: PASS (all %zu elements match 0x%02X)\n",
             Nmemset, (unsigned char)memsetval);
    }
    printf("-------------------\n");
  
    HIP_CHECK(hipFree(A_d));
    HIP_CHECK(hipFree(wid_d));
    HIP_CHECK(hipFree(xcc_d));
    HIP_CHECK(hipFree(xcd_count_d));
    free(A_h);
    free(wid_h);
    free(xcc_h);
    return testResult;
  }

static void printUsage(const char* prog) {
  printf("Usage: %s [options]\n", prog);
  printf("Options:\n");
  printf("  --test-xcd <N>   Run testhipMemsetXcd targeting XCD N\n");
  printf("  -n <N>           Number of elements (default: 4*1024*1024 = 4194304)\n");
  printf("\nIf no options are given, compile-time #ifdef tests are run.\n");
}

int main(int argc, char* argv[]) {
  bool ret;
  printf("Running..\n");

  size_t numElements = 4 * 1024 * 1024;  // default
  bool numElementsSet = false;
  bool ran_cli_test = false;
  int test_xcd_val = -1;  // -1 means not set

  // First pass: parse all options
  for (int i = 1; i < argc; i++) {
    std::string arg = argv[i];
    if (arg == "-n") {
      if (i + 1 >= argc) {
        printf("Error: -n requires a value (e.g. -n 4194304)\n");
        printUsage(argv[0]);
        return 1;
      }
      numElements = (size_t)strtol(argv[i + 1], nullptr, 10);
      numElementsSet = true;
      i++;
    } else if (arg == "--test-xcd") {
      if (i + 1 >= argc) {
        printf("Error: --test-xcd requires a value (e.g. --test-xcd 0)\n");
        printUsage(argv[0]);
        return 1;
      }
      test_xcd_val = (int)strtol(argv[i + 1], nullptr, 10);
      i++;
    } else if (arg == "--help" || arg == "-h") {
      printUsage(argv[0]);
      return 0;
    } else {
      printf("Unknown option: %s\n", arg.c_str());
      printUsage(argv[0]);
      return 1;
    }
  }

  // Run --test-xcd if requested
  if (test_xcd_val >= 0) {
    size_t xcd = (size_t)test_xcd_val;
    printf("Running testhipMemsetXcd with xcd=%zu numElements=%zu\n", xcd, numElements);
    ret = testhipMemsetXcd(0x42, numElements, xcd, 3); assert(ret);
    ran_cli_test = true;
  }

  if (!ran_cli_test) {
#ifdef POISON_INIT_OFFSET_0
    // Fails with offset = 0 consistently
    size_t n0 = numElementsSet ? numElements : 256 * 1024 * 1024;
    ret = testhipMemset(0xa6, n0); assert(ret);
#endif

#ifdef POISON_INIT_OFFSET_3
    // Fails more reliably with offset = 3 with below combinations.
    ret = testhipMemset(0x42, numElements, 3); assert(ret);
#endif

#ifdef POISON_INIT_OFFSET_3_XCD_0
    ret = testhipMemsetXcd(0x42, numElements, 0, 3); assert(ret);
#endif
  }

  printf("\nAll tests PASSED.\n");
  return 0;
}