Skip to content

Instantly share code, notes, and snippets.

@cassc
Last active April 1, 2025 02:01
Show Gist options
  • Save cassc/090f9113c603a7dbe891bfa24b3ddc62 to your computer and use it in GitHub Desktop.
Save cassc/090f9113c603a7dbe891bfa24b3ddc62 to your computer and use it in GitHub Desktop.
cuda: demo no arg constructor in extended struct
/// nvcc -std=c++20 -arch=sm_86 -o demo-no-arg-constructor demo_no_arg_constructor.cu
/// ❯ /opt/cuda/extras/compute-sanitizer/compute-sanitizer \
// --leak-check=full --report-api-errors=all \
// --check-cache-control \
// --check-api-memory-access=yes --check-device-heap=yes \
// --tool memcheck \
// ./demo-no-arg-constructor
#include <cuda_runtime.h>
#include <cstdint>
#include <iostream>
#define CUDA_CHECK(call) do { \
cudaError_t err = call; \
if (err != cudaSuccess) { \
std::cerr << "CUDA call failed: " << #call << "\n" \
<< "Error code: " << err << "\n" \
<< "Error string: " << cudaGetErrorString(err) \
<< " (at " << __FILE__ << ":" << __LINE__ << ")" \
<< std::endl; \
throw std::runtime_error(cudaGetErrorString(err)); \
} \
} while (0)
struct uint256 {
uint32_t words[8];
};
struct evm_word_t : uint256 {
// This default constructor will cause compute-sanitizer error:
// ========= Invalid __global__ write of size 8 bytes
// ========= at test_new_array(unsigned long)+0x110
// ========= by thread (0,0,0) in block (0,0,0)
// ========= Address 0x205dff918 is out of bounds
__host__ __device__ evm_word_t() {}
};
__global__ void test_new_array(size_t N) {
evm_word_t* arr = new evm_word_t[N]; // for large N this can fail with nullptr
if (arr == nullptr) {
printf("Thread %d: new failed\n", threadIdx.x);
return;
}
const int x = blockIdx.x * blockDim.x + threadIdx.x;
arr[0].words[0] = x;
printf("Thread %d: arr[0].words[0] = %u\n", threadIdx.x, arr[0].words[0]);
delete[] arr;
}
int main( int argc, char** argv ) {
test_new_array<<<1, 32>>>(10);
CUDA_CHECK(cudaDeviceSynchronize());
return 0;
}
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment