memset_offset_backup.cpp

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

#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;
    if (i < N) {
        unsigned smid = __smid();
        wid[i] = smid;
        xcc[i] = (smid >> 9) & 0xF;
        dst[i] = value;
        if (dst[ N - i - 1] != value) {
            //printf("Mismatch at i=%zu: expected 0x%02X got 0x%02X\n",
            //       i, value, dst[i]);
            mismatch[i] = true;
            //abort();
        }
    }
}


static bool testhipMemset(char memsetval, size_t numElements, size_t offset = 0) {
  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());

  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]);\
    }
  }


  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));
  free(A_h);
  free(wid_h);
  free(xcc_h);
  return testResult;
}

int main() {
  bool ret;
  printf("Running..\n");


#ifdef POISON_INIT_OFFSET_0
  // Fails with offset = 0 consistently
  ret = testhipMemset(0xa6, 256 * 1024 * 1024); assert(ret);
#endif

#ifdef POISON_INIT_OFFSET_3
  // Fails more reliably with offset = 3 with below combinations.
  ret = testhipMemset(0x42, 4 * 1024 * 1024, 3); assert(ret);
#endif
  printf("\nAll tests PASSED.\n");
  return 0;
}