|
| 1 | +// XFAIL: * |
| 2 | +// REQUIRES: comgr-has-spirv |
| 3 | + |
| 4 | +// COM: Generate a debuginfo SPIR-V file from a HIP kernel |
| 5 | +// RUN: clang -x hip --offload-arch=amdgcnspirv -nogpulib -nogpuinc \ |
| 6 | +// RUN: --no-gpu-bundle-output --offload-device-only -O3 %s -o %t.dbg.spv -g |
| 7 | + |
| 8 | +// COM: Compile debuginfo SPIR-V source to a relocatable |
| 9 | +// RUN: AMD_COMGR_EMIT_VERBOSE_LOGS=1 AMD_COMGR_REDIRECT_LOGS=stdout \ |
| 10 | +// RUN: spirv-to-reloc %t.dbg.spv %t.dbg.o | FileCheck \ |
| 11 | +// RUN: -check-prefix=CHECK-DBG %s |
| 12 | + |
| 13 | +// COM: Check that debuginfo SPIR-V flags are correctly extracted |
| 14 | +// CHECK-DBG: Driver Job Args: {{.*}} "-gheterogeneous-dwarf=diexpression" "-mllvm -amdgpu-spill-cfi-saved-regs" |
| 15 | + |
| 16 | +#include <cstdlib> |
| 17 | + |
| 18 | +#define __constant__ __attribute__((constant)) |
| 19 | +#define __device__ __attribute__((device)) |
| 20 | +#define __global__ __attribute__((global)) |
| 21 | +#define __host__ __attribute__((host)) |
| 22 | +#define __shared__ __attribute__((shared)) |
| 23 | +#define __managed__ __attribute__((managed)) |
| 24 | +#define __launch_bounds__(...) __attribute__((launch_bounds(__VA_ARGS__))) |
| 25 | + |
| 26 | +struct dim3 { |
| 27 | + unsigned x, y, z; |
| 28 | + __host__ __device__ dim3(unsigned x, unsigned y = 1, unsigned z = 1) : x(x), y(y), z(z) {} |
| 29 | +}; |
| 30 | + |
| 31 | +#ifdef __HIP__ |
| 32 | +typedef struct hipStream *hipStream_t; |
| 33 | +typedef enum hipError {} hipError_t; |
| 34 | +int hipConfigureCall(dim3 gridSize, dim3 blockSize, size_t sharedSize = 0, |
| 35 | + hipStream_t stream = 0); |
| 36 | +extern "C" hipError_t __hipPushCallConfiguration(dim3 gridSize, dim3 blockSize, |
| 37 | + size_t sharedSize = 0, |
| 38 | + hipStream_t stream = 0); |
| 39 | +extern "C" hipError_t hipLaunchKernel(const void *func, dim3 gridDim, |
| 40 | + dim3 blockDim, void **args, |
| 41 | + size_t sharedMem, |
| 42 | + hipStream_t stream); |
| 43 | +#endif |
| 44 | + |
| 45 | +__attribute__((device)) |
| 46 | +void clean_value(float* ptr) { *ptr = 0; } |
| 47 | + |
| 48 | +__attribute__((global)) |
| 49 | +void add_value(float* a, float* b, float* res) { |
| 50 | + *res = *a + *b; |
| 51 | + |
| 52 | + clean_value(a); |
| 53 | +} |
0 commit comments