[llvm] Extend llvm objdump fatbin (PR #114834)

Jacob Lambert via llvm-commits llvm-commits at lists.llvm.org
Tue Nov 12 11:04:24 PST 2024


================
@@ -0,0 +1,80 @@
+// RUN: clang++ -x hip --hip-link --offload-arch=gfx1100 --offload-arch=gfx1101 -o %t %s
+// RUN: llvm-objdump %t --offload-fatbin
+// RUN: llvm-objdump %t --offload-fatbin --arch-name=gfx1100
+
+#include <stdio.h>
+
+#include <iostream>
+#include "hip/hip_runtime.h"
+
+__global__ void simpleAdd(uint32_t* A_d, const uint32_t* B_d, size_t N)
+{
+  size_t i = (blockIdx.x * blockDim.x + threadIdx.x);
+  A_d[i] += B_d[i];
+}
+
+int main()
+{
+  int device_count = 0;
+  hipGetDeviceCount(&device_count);
+	
+  std::cout<< "Found " << device_count << " HIP devices." << std::endl;
+  for( int i = 0; i < device_count; i++ )
+  {
+    std::cout << "Device " << i << std::endl;
+    hipDeviceProp_t props;
+    hipGetDeviceProperties(&props, i);
+    std::cout << " Name: " << props.name << std::endl;
+  }
+
+  hipSetDevice(0);
+
+  uint32_t *A_d, *B_d;
+  uint32_t *A_h, *B_h;
+
+  size_t N = 1000;
+  size_t Nbytes = N * sizeof( uint32_t );
+
+  // Allocating host memory
+  A_h = (uint32_t*) malloc(Nbytes);
+  B_h = (uint32_t*) malloc(Nbytes);
+
+  for( size_t i = 0; i < N ; i++ )
+  {
+    A_h[i] = i;
+    B_h[i] = 2*i;
+  }
+
+  // Allocating device memory
+  hipMalloc(&A_d, Nbytes);
+  hipMalloc(&B_d, Nbytes);
+
+  // Copy host to device
+  hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice);
+  hipMemcpy(B_d, B_h, Nbytes, hipMemcpyHostToDevice);
+
+  // launch kernel
+  const unsigned blocks = 512;
+  const unsigned threadsPerBlock = 256;
+  hipLaunchKernelGGL(simpleAdd, dim3(blocks), dim3(threadsPerBlock), 0, 0, A_d, B_d, N);
+
+  // Copy device to host
+  hipMemcpy(A_h, A_d, Nbytes, hipMemcpyDeviceToHost);
+
+  // Verify
+  for( size_t i = 0; i < N; i++ )
+  {
+    std::cout << A_h[i] << " ";
+    uint32_t A_ref = 3*i;
+    if( A_h[i] != A_ref )
+    {
+      std::cout<< "Mismatch occured at " << i << ": " << A_h[i] << " != " << A_ref << std::endl;
----------------
lamb-j wrote:

space after cout?

https://github.com/llvm/llvm-project/pull/114834


More information about the llvm-commits mailing list