AMD hipcc 生成各个gpu 微架构汇编语言代码的方法示例

最新推荐文章于 2024-05-01 21:37:48 发布
最新推荐文章于 2024-05-01 21:37:48 发布
阅读量429 收藏 1
点赞数 15
文章标签： 算法 编译器 高性能计算
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/eloudy/article/details/137178054
版权

1，gpu vectorAdd 示例

为了简化逻辑，故假设 vector 的 size 与运行配置的thread个熟正好一样多，比如都是512之类的.

1.1 源码

vectorAdd.hip

#include <stdio.h>
#include <hip/hip_runtime.h>

__global__ void vectorAdd(const float *A, const float *B, float *C) {
  int i = blockDim.x * blockIdx.x + threadIdx.x;

  C[i] = A[i] + B[i] + 0.0f;
}

int main(void) {
  hipError_t err = hipSuccess;
  int numElements = 512;
  size_t size = numElements * sizeof(float);
  printf("[Vector addition of %d elements]\n", numElements);

  float *h_A = (float *)malloc(size);
  float *h_B = (float *)malloc(size);
  float *h_C = (float *)malloc(size);

  if (h_A == NULL || h_B == NULL || h_C == NULL) {
    fprintf(stderr, "Failed to allocate host vectors!\n");
    exit(EXIT_FAILURE);
  }

  for (int i = 0; i < numElements; ++i) {
    h_A[i] = rand() / (float)RAND_MAX;
    h_B[i] = rand() / (float)RAND_MAX;
  }

  float *d_A = NULL;
  err = hipMalloc((void **)&d_A, size);

  if (err != hipSuccess) {
    fprintf(stderr, "Failed to allocate device vector A (error code %s)!\n",
            hipGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  float *d_B = NULL;
  err = hipMalloc((void **)&d_B, size);

  if (err != hipSuccess) {
    fprintf(stderr, "Failed to allocate device vector B (error code %s)!\n",
            hipGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  float *d_C = NULL;
  err = hipMalloc((void **)&d_C, size);

  if (err != hipSuccess) {
    fprintf(stderr, "Failed to allocate device vector C (error code %s)!\n",
            hipGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  printf("Copy input data from the host memory to the CUDA device\n");
  err = hipMemcpy(d_A, h_A, size, hipMemcpyHostToDevice);

  if (err != hipSuccess) {
    fprintf(stderr,
            "Failed to copy vector A from host to device (error code %s)!\n",
            hipGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  err = hipMemcpy(d_B, h_B, size, hipMemcpyHostToDevice);

  if (err != hipSuccess) {
    fprintf(stderr,
            "Failed to copy vector B from host to device (error code %s)!\n",
            hipGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  int threadsPerBlock = 256;
  int blocksPerGrid = (numElements + threadsPerBlock - 1) / threadsPerBlock;
  printf("CUDA kernel launch with %d blocks of %d threads\n", blocksPerGrid,
         threadsPerBlock);
  vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C);
  err = hipGetLastError();

  if (err != hipSuccess) {
    fprintf(stderr, "Failed to launch vectorAdd kernel (error code %s)!\n",
            hipGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  printf("Copy output data from the CUDA device to the host memory\n");
  err = hipMemcpy(h_C, d_C, size, hipMemcpyDeviceToHost);

  if (err != hipSuccess) {
    fprintf(stderr,
            "Failed to copy vector C from device to host (error code %s)!\n",
            hipGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  for (int i = 0; i < numElements; ++i) {
    if (fabs(h_A[i] + h_B[i] - h_C[i]) > 1e-5) {
      fprintf(stderr, "Result verification failed at element %d!\n", i);
      exit(EXIT_FAILURE);
    }
  }

  printf("Test PASSED\n");

  err = hipFree(d_A);

  if (err != hipSuccess) {
    fprintf(stderr, "Failed to free device vector A (error code %s)!\n",
            hipGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  err = hipFree(d_B);

  if (err != hipSuccess) {
    fprintf(stderr, "Failed to free device vector B (error code %s)!\n",
            hipGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  err = hipFree(d_C);

  if (err != hipSuccess) {
    fprintf(stderr, "Failed to free device vector C (error code %s)!\n",
            hipGetErrorString(err));
    exit(EXIT_FAILURE);
  }

  free(h_A);
  free(h_B);
  free(h_C);

  printf("Done\n");
  return 0;
}

Makefile

all: vectorAdd.hip
	hipcc $< -o vectorAdd
	hipcc -S --cuda-device-only --offload-arch=gfx906 --offload-arch=gfx908 $<

.PHONY: clean
clean:
	rm -rf *.s vectorAdd


# --offload-arch=gfx803 
# --offload-arch=gfx900 
# --offload-arch=gfx906 
# --offload-arch=gfx908 
# --offload-arch=gfx90a 
# --offload-arch=gfx942 
# --offload-arch=gfx1030 
# --offload-arch=gfx1100 
# --offload-arch=gfx1101 
# --offload-arch=gfx1102

1.2  编译运行

make

2，汇编代码

这句命令可以帮助生成汇编代码：

hipcc -S --cuda-device-only --offload-arch=gfx906 --offload-arch=gfx908 $<

offload-arch 帮助指定 gpu 的微架构，

vectorAdd-hip-amdgcn-amd-amdhsa-gfx906.s :

	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx906"
	.protected	_Z9vectorAddPKfS0_Pf    ; -- Begin function _Z9vectorAddPKfS0_Pf
	.globl	_Z9vectorAddPKfS0_Pf
	.p2align	8
	.type	_Z9vectorAddPKfS0_Pf,@function
_Z9vectorAddPKfS0_Pf:                   ; @_Z9vectorAddPKfS0_Pf
; %bb.0:
	s_load_dword s7, s[4:5], 0x24
	s_load_dwordx4 s[0:3], s[4:5], 0x0
	s_load_dwordx2 s[8:9], s[4:5], 0x10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s4, s7, 0xffff
	s_mul_i32 s6, s6, s4
	v_add_u32_e32 v0, s6, v0
	v_ashrrev_i32_e32 v1, 31, v0
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	v_mov_b32_e32 v3, s1
	v_add_co_u32_e32 v2, vcc, s0, v0
	v_addc_co_u32_e32 v3, vcc, v3, v1, vcc
	global_load_dword v4, v[2:3], off
	v_mov_b32_e32 v3, s3
	v_add_co_u32_e32 v2, vcc, s2, v0
	v_addc_co_u32_e32 v3, vcc, v3, v1, vcc
	global_load_dword v2, v[2:3], off
	v_mov_b32_e32 v3, s9
	v_add_co_u32_e32 v0, vcc, s8, v0
	v_addc_co_u32_e32 v1, vcc, v3, v1, vcc
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v2, v4, v2
	v_add_f32_e32 v2, 0, v2
	global_store_dword v[0:1], v2, off
	s_endpgm
	.section	.rodata,#alloc
	.p2align	6, 0x0
	.amdhsa_kernel _Z9vectorAddPKfS0_Pf
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.amdhsa_user_sgpr_count 6
		.amdhsa_user_sgpr_private_segment_buffer 1
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_flat_scratch_init 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_system_sgpr_private_segment_wavefront_offset 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 5
		.amdhsa_next_free_sgpr 10
		.amdhsa_reserve_flat_scratch 0
		.amdhsa_reserve_xnack_mask 1
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z9vectorAddPKfS0_Pf, .Lfunc_end0-_Z9vectorAddPKfS0_Pf
                                        ; -- End function
	.section	.AMDGPU.csdata
; Kernel info:
; codeLenInByte = 132
; NumSgprs: 14
; NumVgprs: 5
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 14
; NumVGPRsForWavesPerEU: 5
; Occupancy: 8
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 6
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.protected	_ZN17__HIP_CoordinatesI14__HIP_BlockDimE1xE ; @_ZN17__HIP_CoordinatesI14__HIP_BlockDimE1xE
	.type	_ZN17__HIP_CoordinatesI14__HIP_BlockDimE1xE,@object
	.section	.rodata._ZN17__HIP_CoordinatesI14__HIP_BlockDimE1xE,#alloc
	.weak	_ZN17__HIP_CoordinatesI14__HIP_BlockDimE1xE
_ZN17__HIP_CoordinatesI14__HIP_BlockDimE1xE:
	.zero	1
	.size	_ZN17__HIP_CoordinatesI14__HIP_BlockDimE1xE, 1

	.protected	_ZN17__HIP_CoordinatesI14__HIP_BlockIdxE1xE ; @_ZN17__HIP_CoordinatesI14__HIP_BlockIdxE1xE
	.type	_ZN17__HIP_CoordinatesI14__HIP_BlockIdxE1xE,@object
	.section	.rodata._ZN17__HIP_CoordinatesI14__HIP_BlockIdxE1xE,#alloc
	.weak	_ZN17__HIP_CoordinatesI14__HIP_BlockIdxE1xE
_ZN17__HIP_CoordinatesI14__HIP_BlockIdxE1xE:
	.zero	1
	.size	_ZN17__HIP_CoordinatesI14__HIP_BlockIdxE1xE, 1

	.protected	_ZN17__HIP_CoordinatesI15__HIP_ThreadIdxE1xE ; @_ZN17__HIP_CoordinatesI15__HIP_ThreadIdxE1xE
	.type	_ZN17__HIP_CoordinatesI15__HIP_ThreadIdxE1xE,@object
	.section	.rodata._ZN17__HIP_CoordinatesI15__HIP_ThreadIdxE1xE,#alloc
	.weak	_ZN17__HIP_CoordinatesI15__HIP_ThreadIdxE1xE
_ZN17__HIP_CoordinatesI15__HIP_ThreadIdxE1xE:
	.zero	1
	.size	_ZN17__HIP_CoordinatesI15__HIP_ThreadIdxE1xE, 1

	.ident	"AMD clang version 17.0.0 (https://github.com/RadeonOpenCompute/llvm-project roc-6.0.2 24012 af27734ed982b52a9f1be0f035ac91726fc697e4)"
	.section	".note.GNU-stack"
	.addrsig
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z9vectorAddPKfS0_Pf
    .private_segment_fixed_size: 0
    .sgpr_count:     14
    .sgpr_spill_count: 0
    .symbol:         _Z9vectorAddPKfS0_Pf.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     5
    .vgpr_spill_count: 0
    .wavefront_size: 64
amdhsa.target:   amdgcn-amd-amdhsa--gfx906
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata

vectorAdd-hip-amdgcn-amd-amdhsa-gfx908.s :


	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx908"
	.protected	_Z9vectorAddPKfS0_Pf    ; -- Begin function _Z9vectorAddPKfS0_Pf
	.globl	_Z9vectorAddPKfS0_Pf
	.p2align	8
	.type	_Z9vectorAddPKfS0_Pf,@function
_Z9vectorAddPKfS0_Pf:                   ; @_Z9vectorAddPKfS0_Pf
; %bb.0:
	s_load_dword s7, s[4:5], 0x24
	s_load_dwordx4 s[0:3], s[4:5], 0x0
	s_load_dwordx2 s[8:9], s[4:5], 0x10
	s_waitcnt lgkmcnt(0)
	s_and_b32 s4, s7, 0xffff
	s_mul_i32 s6, s6, s4
	v_add_u32_e32 v0, s6, v0
	v_ashrrev_i32_e32 v1, 31, v0
	v_lshlrev_b64 v[0:1], 2, v[0:1]
	v_mov_b32_e32 v3, s1
	v_add_co_u32_e32 v2, vcc, s0, v0
	v_addc_co_u32_e32 v3, vcc, v3, v1, vcc
	global_load_dword v4, v[2:3], off
	v_mov_b32_e32 v3, s3
	v_add_co_u32_e32 v2, vcc, s2, v0
	v_addc_co_u32_e32 v3, vcc, v3, v1, vcc
	global_load_dword v2, v[2:3], off
	v_mov_b32_e32 v3, s9
	v_add_co_u32_e32 v0, vcc, s8, v0
	v_addc_co_u32_e32 v1, vcc, v3, v1, vcc
	s_waitcnt vmcnt(0)
	v_add_f32_e32 v2, v4, v2
	v_add_f32_e32 v2, 0, v2
	global_store_dword v[0:1], v2, off
	s_endpgm
	.section	.rodata,#alloc
	.p2align	6, 0x0
	.amdhsa_kernel _Z9vectorAddPKfS0_Pf
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 0
		.amdhsa_kernarg_size 280
		.amdhsa_user_sgpr_count 6
		.amdhsa_user_sgpr_private_segment_buffer 1
		.amdhsa_user_sgpr_dispatch_ptr 0
		.amdhsa_user_sgpr_queue_ptr 0
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 0
		.amdhsa_user_sgpr_flat_scratch_init 0
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_uses_dynamic_stack 0
		.amdhsa_system_sgpr_private_segment_wavefront_offset 0
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 0
		.amdhsa_system_sgpr_workgroup_id_z 0
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 0
		.amdhsa_next_free_vgpr 5
		.amdhsa_next_free_sgpr 10
		.amdhsa_reserve_flat_scratch 0
		.amdhsa_reserve_xnack_mask 1
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z9vectorAddPKfS0_Pf, .Lfunc_end0-_Z9vectorAddPKfS0_Pf
                                        ; -- End function
	.section	.AMDGPU.csdata
; Kernel info:
; codeLenInByte = 132
; NumSgprs: 14
; NumVgprs: 5
; NumAgprs: 0
; TotalNumVgprs: 5
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 14
; NumVGPRsForWavesPerEU: 5
; Occupancy: 8
; WaveLimiterHint : 1
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 6
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
	.protected	_ZN17__HIP_CoordinatesI14__HIP_BlockDimE1xE ; @_ZN17__HIP_CoordinatesI14__HIP_BlockDimE1xE
	.type	_ZN17__HIP_CoordinatesI14__HIP_BlockDimE1xE,@object
	.section	.rodata._ZN17__HIP_CoordinatesI14__HIP_BlockDimE1xE,#alloc
	.weak	_ZN17__HIP_CoordinatesI14__HIP_BlockDimE1xE
_ZN17__HIP_CoordinatesI14__HIP_BlockDimE1xE:
	.zero	1
	.size	_ZN17__HIP_CoordinatesI14__HIP_BlockDimE1xE, 1

	.protected	_ZN17__HIP_CoordinatesI14__HIP_BlockIdxE1xE ; @_ZN17__HIP_CoordinatesI14__HIP_BlockIdxE1xE
	.type	_ZN17__HIP_CoordinatesI14__HIP_BlockIdxE1xE,@object
	.section	.rodata._ZN17__HIP_CoordinatesI14__HIP_BlockIdxE1xE,#alloc
	.weak	_ZN17__HIP_CoordinatesI14__HIP_BlockIdxE1xE
_ZN17__HIP_CoordinatesI14__HIP_BlockIdxE1xE:
	.zero	1
	.size	_ZN17__HIP_CoordinatesI14__HIP_BlockIdxE1xE, 1

	.protected	_ZN17__HIP_CoordinatesI15__HIP_ThreadIdxE1xE ; @_ZN17__HIP_CoordinatesI15__HIP_ThreadIdxE1xE
	.type	_ZN17__HIP_CoordinatesI15__HIP_ThreadIdxE1xE,@object
	.section	.rodata._ZN17__HIP_CoordinatesI15__HIP_ThreadIdxE1xE,#alloc
	.weak	_ZN17__HIP_CoordinatesI15__HIP_ThreadIdxE1xE
_ZN17__HIP_CoordinatesI15__HIP_ThreadIdxE1xE:
	.zero	1
	.size	_ZN17__HIP_CoordinatesI15__HIP_ThreadIdxE1xE, 1

	.ident	"AMD clang version 17.0.0 (https://github.com/RadeonOpenCompute/llvm-project roc-6.0.2 24012 af27734ed982b52a9f1be0f035ac91726fc697e4)"
	.section	".note.GNU-stack"
	.addrsig
	.amdgpu_metadata
---
amdhsa.kernels:
  - .agpr_count:     0
    .args:
      - .address_space:  global
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .offset:         16
        .size:           8
        .value_kind:     global_buffer
      - .offset:         24
        .size:           4
        .value_kind:     hidden_block_count_x
      - .offset:         28
        .size:           4
        .value_kind:     hidden_block_count_y
      - .offset:         32
        .size:           4
        .value_kind:     hidden_block_count_z
      - .offset:         36
        .size:           2
        .value_kind:     hidden_group_size_x
      - .offset:         38
        .size:           2
        .value_kind:     hidden_group_size_y
      - .offset:         40
        .size:           2
        .value_kind:     hidden_group_size_z
      - .offset:         42
        .size:           2
        .value_kind:     hidden_remainder_x
      - .offset:         44
        .size:           2
        .value_kind:     hidden_remainder_y
      - .offset:         46
        .size:           2
        .value_kind:     hidden_remainder_z
      - .offset:         64
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         72
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         80
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         88
        .size:           2
        .value_kind:     hidden_grid_dims
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 280
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z9vectorAddPKfS0_Pf
    .private_segment_fixed_size: 0
    .sgpr_count:     14
    .sgpr_spill_count: 0
    .symbol:         _Z9vectorAddPKfS0_Pf.kd
    .uniform_work_group_size: 1
    .uses_dynamic_stack: false
    .vgpr_count:     5
    .vgpr_spill_count: 0
    .wavefront_size: 64
amdhsa.target:   amdgcn-amd-amdhsa--gfx908
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata

对比可以发现，两个kernel的 asm 内容差不多，说明两个微架构差不多。

Eloudy
关注
  • 15
    点赞
  • 踩
  • 1
    收藏
    觉得还不错? 一键收藏
  • 0
    评论
Langchain+本地大语言模型进行数据库操作的实战代码
herosunly的博客
06-23 8万+
本文讲解了Langchain+本地大语言模型进行数据库操作的实战代码，希望能对尝试使用开源大语言模型进行SQL操作的同学们有所帮助。 文章目录 1. 前言 2. 代码思路剖析 3. 实战代码
transformers加载模型时自动选择空闲的单个GPU的实战代码
热门推荐
weixin_43178406的博客
07-20 6万+
本文主要介绍了transformers加载模型时自动选择空闲GPU的实战代码，希望能对使用huggingface transformers的同学们有所帮助。 文章目录 1. 背景描述 2. 实战 2.1 安装Python库 2.2 实战代码
GPU指令下发方式
kayshi的博客
10-24 885
描述GPU指令下发方式
x86、amd、arm和GPU
左直拳的马桶_日用桶
04-10 1万+
参考资料： Intel和AMD 与 x86，ARM，MIPS有什么区别？
windows10系统查看NVIDIA显卡GPU使用情况——解决方法
黄笳倞的博客
02-14 4万+
目录 一、前言 二、解决方法 一、前言 1.随着深度学习的快速发展，GPU的优劣也决定模型训练速度的快慢，而了解GPU的使用情况可以让我们合理的分配GPU。 2.现在普遍的深度学习的GPU采用的都是NVIDIA显卡，而企业也提供了专门的方式供我们去查看GPU的使用情况。对于一些windows10系统的用户，他们可以在任务管理器的性能界面查看到GPU的使用情况，但是有些用户就不可以，例如博...
clrx_OpencL_GPU_汇编_amd_
10-01
AMD GPU的汇编编程工具，可以直接使用AMD底层API进行代码编写
three.js利用gpu选取物体并计算交点位置的方法示例
12-13
使用three.js自带的光线投射器(Raycaster)选取物体非常简单，代码如下所示： var raycaster = new THREE.Raycaster(); var mouse = new THREE.Vector2(); function onMouseMove(event) { // 计算鼠标所在位置的...
linux-firmware-amdgpu:修复W
03-05
lib /固件/ amdgpu 修复W: Possible missing firmware /lib/firmware/amdgpu/ 修复步骤 打开终端，逐步复制并插入这些逗号。 第1步git clone "https://github.com/Ritwikrajsingh/linux-firmware-amdgpu.git" ...
GPU高性能编程CUDA实战—示例代码
10-28
该代码集是GPU高性能编程CUDA实战一书的完全代码，如有问题请联系博主。
NVIDIA GPU图灵架构详解
01-10
NVIDIA GPU图灵架构详解，NVIDIA GPU图灵架构详解，NVIDIA GPU图灵架构详解，NVIDIA GPU图灵架构详解，NVIDIA GPU图灵架构详解，NVIDIA GPU图灵架构详解。
Rapidly exploring Random Trees(RRT)类算法
qq_35635374的博客
04-28 168
认知有限，望大家多多包涵，有什么问题也希望能够与大家多交流，共同成长！本文先对Rapidly exploring Random Trees(RRT)类算法做个简单的介绍，具体内容后续再更，其他模块可以参考去我其他文章提示：以下是本篇文章正文内容（1）优点快速随机树算法是一种较为迅捷搜索的方式，增量式的构建树其优点明显，主要体现在针对性的找一条从起点到终点的路径其较强较快的搜索能力，【防盗标记–盒子君hzj】并且运用快速拓展随机树的路径规划对路径运行环境没有限制要求。（2）缺点。
算法训练营第十一天 | LeetCode 20 有效的括号、LeetCode 1047 删除字符串中的所有相邻重复项
qq_63149342的博客
04-27 134
注意用的是字符串数组作为输入，而我们要用int型进行计算，这时候就要用到一个string转int类型的库函数stoi函数，aoti是字符数组转int类型。一个是最后循环结束后要判断栈内存是否为空，一个是出栈时要判断栈是否为空，总的来说是考验对于栈内存空间的掌控吧。这一题考栈的性质，相邻字符用未进栈的元素和进栈的元素进行模拟即可，相同且栈不为空即出栈，否则进栈，之后将栈内元素全部导入一个字符串，还要注意要写一个反转函数翻转下即可。LeetCode 1047 删除字符串中的所有相邻重复项。
c++容器与算法概述
最新发布
北风
05-01 309
简单讲述输入输出运算符重载 以及容器的简单使用及接口哦苦
代码随想录算法训练营第三十八天
weixin_41902984的博客
04-28 206
又又又欠两天债，不过回家了，一切可以按部就班，明天慢慢补吧，今天先做一课的内容。然后就是如果一次可以上m个阶梯的话应该怎么做，先把伪代码放上来。其实就是斐波那契数，不过需要理解一下过程。还好还好，动态规划第一天不太难。746. 使用最小花费爬楼梯。509. 斐波那契数。
Golang | Leetcode Golang题解之第56题合并区间
weixin_66442839的博客
04-28 479
Golang | Leetcode Golang题解之第56题合并区间
4/26发布发布：缺了好几次的作业，矩形法+二分法求下面方程根+顺序查找n+程序填空，补一下还有八九没做，炸8412 字不是干的，哈哈哈
2201_75637762的博客
04-26 1246
eps：用于定义逼近的精度，即当x1和x2的差小于eps时，认为找到了根。x1和x2：用于二分法的初始区间，这里选择的是[-10, 10]。y1和y2：分别是x1和x2处函数的值。这个程序通过二分法在一个给定的区间内寻找三次方程的根。它首先检查函数在区间端点的值的符号，然后使用二分法逐步缩小区间，直到找到满足精度要求的根或确定区间内没有根为止。okll测试数据：3 4 7 12 24 78 9 15 80 45第一次运行程序输入测试数据 12。
c入门基础题（3）
梦昼初
04-26 901
1、找出数组中最大元素的位置（下标值） 输入 多组测试，每组先输入一个不大于10的整数n 然后是n个整数 输出 输出这n个整数中最大的元素及下标值 样例输入 4 1 4 5 6 样例输出 6 3 #include<stdio.h> int main(){ int a[20]; int n,x=0,y,z,s,m,f; int i,j; while(scanf("%d",&n)!=EOF){ for(i=1;i<=n;i++){ scanf("%d",&y)
【C++中的模板】
Z15255808023的博客
04-27 739
本片内容主要是C++中模板的有关知识，主要讲述了模板及其实用和注意事项，请耐心观看。相必我们大家都听说过泛型编程这个概念：编写与类型无关的通用代码，是代码复用的一种手段。模板是泛型编程的基础。简单的说就是编写一种可以在多种数据类型上都能用的代码。比如说，你写一次代码，就可以在很多不同的数据类型上使用，而不需要为每种数据类型都重新写一遍相同逻辑的代码。它能让编程变得更简单、更灵活。
测试tensorflow cpu和gpu计算性能的示例代码
05-31
以下是测试TensorFlow CPU和GPU计算性能的示例代码： ```python import tensorflow as tf import timeit # 创建一个随机矩阵 random_matrix = tf.random.normal([1000, 1000]) # 测试CPU性能 with tf.device('/...

“相关推荐”对你有帮助么？

  • 非常没帮助
  • 没帮助
  • 一般
  • 有帮助
  • 非常有帮助
提交
评论
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值