[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