Если вы используете графический процессор Nvidia, вы можете просмотреть код сборки PTX . PTX - это всего лишь псевдо-сборка, своего рода промежуточный между OpenCL и двоичным кодом, который фактически выполняется на GPU.
Вот как вы попали к нему из OpenCL:
Context context(device);
queue = CommandQueue(context, device); // queue to push commands for the device
Program::Sources source;
string kernel_code = opencl_code_settings(N,M)+opencl_code();
source.push_back({ kernel_code.c_str(), kernel_code.length() });
Program program(context, source);
if(program.build("-cl-fast-relaxed-math")) return false; // compile OpenCL code, return false if there is an error
const string ptx_code = program.getInfo<CL_PROGRAM_BINARIES>()[0]; // generate assembly (ptx) for OpenCL code
Строка ptx_code - это то, что вы ищете. Вот небольшой пример ядра:
kernel void benchmark_1(global float* data) {
const uint n = get_global_id(0);
#pragma unroll
for(uint i=0; i<def_M; i++) data[i*def_N+n] = 0.0f;
}
Вот как выглядит код PTX для этого ядра:
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: UNKNOWN
// Driver
// Based on LLVM 3.4svn
//
.version 6.2
.target sm_61, texmode_independent
.address_size 64
// .globl benchmark_1
.entry benchmark_1(
.param .u64 .ptr .global .align 4 benchmark_1_param_0
)
{
.reg .b32 %r<23>;
.reg .b64 %rd<34>;
ld.param.u64 %rd1, [benchmark_1_param_0];
mov.b32 %r1, %envreg3;
mov.u32 %r2, %ntid.x;
mov.u32 %r3, %ctaid.x;
mad.lo.s32 %r4, %r3, %r2, %r1;
mov.u32 %r5, %tid.x;
add.s32 %r6, %r4, %r5;
mul.wide.u32 %rd2, %r6, 4;
add.s64 %rd3, %rd1, %rd2;
mov.u32 %r7, 0;
st.global.u32 [%rd3], %r7;
add.s32 %r8, %r6, 15728640;
mul.wide.u32 %rd4, %r8, 4;
add.s64 %rd5, %rd1, %rd4;
st.global.u32 [%rd5], %r7;
add.s32 %r9, %r6, 31457280;
mul.wide.u32 %rd6, %r9, 4;
add.s64 %rd7, %rd1, %rd6;
st.global.u32 [%rd7], %r7;
add.s32 %r10, %r6, 47185920;
mul.wide.u32 %rd8, %r10, 4;
add.s64 %rd9, %rd1, %rd8;
st.global.u32 [%rd9], %r7;
add.s32 %r11, %r6, 62914560;
mul.wide.u32 %rd10, %r11, 4;
add.s64 %rd11, %rd1, %rd10;
st.global.u32 [%rd11], %r7;
add.s32 %r12, %r6, 78643200;
mul.wide.u32 %rd12, %r12, 4;
add.s64 %rd13, %rd1, %rd12;
st.global.u32 [%rd13], %r7;
add.s32 %r13, %r6, 94371840;
mul.wide.u32 %rd14, %r13, 4;
add.s64 %rd15, %rd1, %rd14;
st.global.u32 [%rd15], %r7;
add.s32 %r14, %r6, 110100480;
mul.wide.u32 %rd16, %r14, 4;
add.s64 %rd17, %rd1, %rd16;
st.global.u32 [%rd17], %r7;
add.s32 %r15, %r6, 125829120;
mul.wide.u32 %rd18, %r15, 4;
add.s64 %rd19, %rd1, %rd18;
st.global.u32 [%rd19], %r7;
add.s32 %r16, %r6, 141557760;
mul.wide.u32 %rd20, %r16, 4;
add.s64 %rd21, %rd1, %rd20;
st.global.u32 [%rd21], %r7;
add.s32 %r17, %r6, 157286400;
mul.wide.u32 %rd22, %r17, 4;
add.s64 %rd23, %rd1, %rd22;
st.global.u32 [%rd23], %r7;
add.s32 %r18, %r6, 173015040;
mul.wide.u32 %rd24, %r18, 4;
add.s64 %rd25, %rd1, %rd24;
st.global.u32 [%rd25], %r7;
add.s32 %r19, %r6, 188743680;
mul.wide.u32 %rd26, %r19, 4;
add.s64 %rd27, %rd1, %rd26;
st.global.u32 [%rd27], %r7;
add.s32 %r20, %r6, 204472320;
mul.wide.u32 %rd28, %r20, 4;
add.s64 %rd29, %rd1, %rd28;
st.global.u32 [%rd29], %r7;
add.s32 %r21, %r6, 220200960;
mul.wide.u32 %rd30, %r21, 4;
add.s64 %rd31, %rd1, %rd30;
st.global.u32 [%rd31], %r7;
add.s32 %r22, %r6, 235929600;
mul.wide.u32 %rd32, %r22, 4;
add.s64 %rd33, %rd1, %rd32;
st.global.u32 [%rd33], %r7;
ret;
}
Из кода PTX вы можете, например, сосчитать FLOP и передачи памяти, чтобы проверить, насколько эффективно код выполняется через модель линии крыши.