COMPILER]: COMPILE ERROR: No functions with semantic types found

[790989228]

this is my test.cu code

#include <crt/host_defines.h>
#include <optix_device.h>
#include "Interaction.h"
#include "LaunchParam.h"
extern "C" __constant__ OptixLaunchParams optixLaunchParams;

extern "C" __global__ void __closesthit__bsdf_radiance()
{
    const MeshSBTData &sbtData
      = *(const MeshSBTData*)optixGetSbtDataPointer();
    PerIntersectionData &prd = *getPRD<PerIntersectionData>();
    const int primID = optixGetPrimitiveIndex();
    const int3 index  = sbtData.index_buffer[primID];
    const float u = optixGetTriangleBarycentrics().x;
    const float v = optixGetTriangleBarycentrics().y;
    const float3 &A     = sbtData.vertex_buffer[index.x];
    const float3 &B     = sbtData.vertex_buffer[index.y];
    const float3 &C     = sbtData.vertex_buffer[index.z];
    float3 Ng = cross(B-A,C-A);
    float3 geometric_normal = (sbtData.normal_buffer)
      ? ((1.f-u-v) * sbtData.normal_buffer[index.x]
         +       u * sbtData.normal_buffer[index.y]
         +       v * sbtData.normal_buffer[index.z])
      : Ng;
    float3 ray_origin = optixGetWorldRayOrigin();
    float3 ray_dir = optixGetWorldRayDirection();
    float t_hit = optixGetRayTmax();
	
    // float3 world_shading_normal = normalize(optixTransformNormalFromObjectToWorldSpace(shading_normal));
    float3 world_geometric_normal = normalize(optixTransformNormalFromObjectToWorldSpace(-geometric_normal));
    float3 world_shading_normal = world_geometric_normal;//todo 暂时先用world_geometric_normal
    float3 ffnormal = faceforward(world_shading_normal, -ray_dir, world_geometric_normal);

    float3 hitpoint = ray_origin + t_hit * ray_dir;
}

and ptx code：

//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-34431801
// Cuda compilation tools, release 12.6, V12.6.20
// Based on NVVM 7.0.1
//

.version 8.5
.target sm_50
.address_size 64

	// .globl	__closesthit__bsdf_radiance
.global .align 4 .b8 __nv_static_28__e1187107_7_test_cu_6c9dc3b4_identity[64] = {0, 0, 128, 63, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 128, 63, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 128, 63, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 128, 63};
.const .align 4 .f32 __nv_static_28__e1187107_7_test_cu_6c9dc3b4_Pi = 0f40490FDB;
.const .align 4 .f32 __nv_static_28__e1187107_7_test_cu_6c9dc3b4_InvPi = 0f3EA2F983;

.visible .entry __closesthit__bsdf_radiance()
{
	.reg .pred 	%p<12>;
	.reg .f32 	%f<44>;
	.reg .b32 	%r<163>;
	.reg .b64 	%rd<132>;
	.loc	1 7 0


	.loc	1 11 5
	.loc	2 61 2, function_name $L__info_string0, inlined_at 1 11 5
	.loc	3 1015 5, function_name $L__info_string1, inlined_at 2 61 2
	mov.u32 	%r6, 0;
	// begin inline asm
	call (%r5), _optix_get_payload, (%r6);
	// end inline asm
	.loc	2 62 2, function_name $L__info_string0, inlined_at 1 11 5
	.loc	3 1022 5, function_name $L__info_string2, inlined_at 2 62 2
	mov.u32 	%r8, 1;
	// begin inline asm
	call (%r7), _optix_get_payload, (%r8);
	// end inline asm
	.loc	1 30 5
	.loc	3 1631 5, function_name $L__info_string3, inlined_at 1 30 5
	.loc	3 1642 5, function_name $L__info_string4, inlined_at 3 1631 5
	// begin inline asm
	call (%r10), _optix_get_transform_list_size, ();
	// end inline asm
	.loc	3 1631 5, function_name $L__info_string3, inlined_at 1 30 5
	setp.eq.s32 	%p1, %r10, 0;
	@%p1 bra 	$L__BB0_15;

	.loc	3 1635 5, function_name $L__info_string3, inlined_at 1 30 5
	.loc	4 340 5, function_name $L__info_string5, inlined_at 3 1635 5
	.loc	3 1642 5, function_name $L__info_string4, inlined_at 4 340 5
	// begin inline asm
	call (%r11), _optix_get_transform_list_size, ();
	// end inline asm
	.loc	4 341 5, function_name $L__info_string5, inlined_at 3 1635 5
	.loc	3 1301 5, function_name $L__info_string6, inlined_at 4 341 5
	// begin inline asm
	call (%f13), _optix_get_ray_time, ();
	// end inline asm
	.loc	4 344 5, function_name $L__info_string5, inlined_at 3 1635 5
	setp.eq.s32 	%p2, %r11, 0;
	@%p2 bra 	$L__BB0_15;

	.loc	4 0 5
	mov.u32 	%r162, 0;

$L__BB0_3:
	.pragma "nounroll";
	.loc	4 346 9, function_name $L__info_string5, inlined_at 3 1635 5
	.loc	3 1649 5, function_name $L__info_string7, inlined_at 4 346 9
	// begin inline asm
	call (%rd12), _optix_get_transform_list_handle, (%r162);
	// end inline asm
	.loc	4 349 9, function_name $L__info_string5, inlined_at 3 1635 5
	.loc	4 292 5, function_name $L__info_string8, inlined_at 4 349 9
	.loc	3 1656 5, function_name $L__info_string9, inlined_at 4 292 5
	// begin inline asm
	call (%r14), _optix_get_transform_type_from_handle, (%rd12);
	// end inline asm
	.loc	4 294 5, function_name $L__info_string8, inlined_at 4 349 9
	or.b32  	%r15, %r14, 1;
	setp.eq.s32 	%p3, %r15, 3;
	@%p3 bra 	$L__BB0_9;
	bra.uni 	$L__BB0_4;

$L__BB0_9:
	.loc	4 296 9, function_name $L__info_string8, inlined_at 4 349 9
	setp.eq.s32 	%p6, %r14, 2;
	@%p6 bra 	$L__BB0_12;
	bra.uni 	$L__BB0_10;

$L__BB0_12:
	.loc	4 298 13, function_name $L__info_string8, inlined_at 4 349 9
	.loc	3 1677 5, function_name $L__info_string10, inlined_at 4 298 13
	// begin inline asm
	call (%rd84), _optix_get_matrix_motion_transform_from_handle, (%rd12);
	// end inline asm
	.loc	4 299 13, function_name $L__info_string8, inlined_at 4 349 9
	.loc	4 249 5, function_name $L__info_string11, inlined_at 4 299 13
	.loc	4 66 9, function_name $L__info_string12, inlined_at 4 249 5
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd86, %rd84;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r103,%r104,%r105,%r106}, [%rd86];
	// end inline asm
	.loc	4 66 9, function_name $L__info_string12, inlined_at 4 249 5
	add.s64 	%rd90, %rd84, 16;
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd89, %rd90;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r107,%r108,%r109,%r110}, [%rd89];
	// end inline asm
	.loc	4 66 9, function_name $L__info_string12, inlined_at 4 249 5
	add.s64 	%rd93, %rd84, 32;
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd92, %rd93;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r111,%r112,%r113,%r114}, [%rd92];
	// end inline asm
	.loc	4 66 9, function_name $L__info_string12, inlined_at 4 249 5
	add.s64 	%rd96, %rd84, 48;
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd95, %rd96;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r115,%r116,%r117,%r118}, [%rd95];
	// end inline asm
	.loc	4 66 9, function_name $L__info_string12, inlined_at 4 249 5
	add.s64 	%rd99, %rd84, 64;
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd98, %rd99;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r119,%r120,%r121,%r122}, [%rd98];
	// end inline asm
	.loc	4 66 9, function_name $L__info_string12, inlined_at 4 249 5
	add.s64 	%rd102, %rd84, 80;
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd101, %rd102;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r123,%r124,%r125,%r126}, [%rd101];
	// end inline asm
	.loc	4 66 9, function_name $L__info_string12, inlined_at 4 249 5
	add.s64 	%rd105, %rd84, 96;
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd104, %rd105;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r127,%r128,%r129,%r130}, [%rd104];
	// end inline asm
	.loc	4 66 9, function_name $L__info_string12, inlined_at 4 249 5
	add.s64 	%rd108, %rd84, 112;
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd107, %rd108;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r131,%r132,%r133,%r134}, [%rd107];
	// end inline asm
	.loc	4 67 5, function_name $L__info_string12, inlined_at 4 249 5
	mov.b32 	%f29, %r106;
	mov.b32 	%f30, %r107;
	.loc	4 249 5, function_name $L__info_string11, inlined_at 4 299 13
	.loc	4 221 5, function_name $L__info_string14, inlined_at 4 249 5
	and.b32  	%r147, %r105, 65535;
	add.s32 	%r148, %r147, -1;
	cvt.rn.f32.s32 	%f31, %r148;
	.loc	4 227 5, function_name $L__info_string14, inlined_at 4 249 5
	sub.ftz.f32 	%f32, %f13, %f29;
	sub.ftz.f32 	%f33, %f30, %f29;
	div.approx.ftz.f32 	%f34, %f32, %f33;
	mul.ftz.f32 	%f35, %f34, %f31;
	.loc	4 227 5, function_name $L__info_string14, inlined_at 4 249 5
	.loc	5 979 3, function_name $L__info_string15, inlined_at 4 227 5
	min.ftz.f32 	%f36, %f31, %f35;
	.loc	4 227 5, function_name $L__info_string14, inlined_at 4 249 5
	.loc	5 1107 3, function_name $L__info_string16, inlined_at 4 227 5
	mov.f32 	%f37, 0f00000000;
	max.ftz.f32 	%f38, %f37, %f36;
	.loc	4 230 5, function_name $L__info_string14, inlined_at 4 249 5
	setp.num.ftz.f32 	%p9, %f38, %f38;
	selp.f32 	%f39, %f38, 0f00000000, %p9;
	.loc	4 233 5, function_name $L__info_string14, inlined_at 4 249 5
	cvt.rmi.ftz.f32.f32 	%f40, %f39;
	add.ftz.f32 	%f41, %f31, 0fBF800000;
	min.ftz.f32 	%f42, %f40, %f41;
	.loc	4 235 5, function_name $L__info_string14, inlined_at 4 249 5
	sub.ftz.f32 	%f43, %f39, %f42;
	.loc	4 236 5, function_name $L__info_string14, inlined_at 4 249 5
	cvt.rzi.ftz.s32.f32 	%r149, %f42;
	.loc	4 252 5, function_name $L__info_string11, inlined_at 4 299 13
	cvt.s64.s32 	%rd10, %r149;
	mul.wide.s32 	%rd119, %r149, 48;
	add.s64 	%rd111, %rd93, %rd119;
	.loc	4 255 5, function_name $L__info_string11, inlined_at 4 299 13
	.loc	4 174 5, function_name $L__info_string17, inlined_at 4 255 5
	.loc	4 66 9, function_name $L__info_string18, inlined_at 4 174 5
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd110, %rd111;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r135,%r136,%r137,%r138}, [%rd110];
	// end inline asm
	.loc	4 175 5, function_name $L__info_string17, inlined_at 4 255 5
	add.s64 	%rd114, %rd111, 16;
	.loc	4 175 5, function_name $L__info_string17, inlined_at 4 255 5
	.loc	4 66 9, function_name $L__info_string18, inlined_at 4 175 5
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd113, %rd114;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r139,%r140,%r141,%r142}, [%rd113];
	// end inline asm
	.loc	4 176 5, function_name $L__info_string17, inlined_at 4 255 5
	add.s64 	%rd117, %rd111, 32;
	.loc	4 176 5, function_name $L__info_string17, inlined_at 4 255 5
	.loc	4 66 9, function_name $L__info_string18, inlined_at 4 176 5
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd116, %rd117;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r143,%r144,%r145,%r146}, [%rd116];
	// end inline asm
	.loc	4 179 5, function_name $L__info_string17, inlined_at 4 255 5
	setp.leu.ftz.f32 	%p10, %f43, 0f00000000;
	@%p10 bra 	$L__BB0_14;

	.loc	4 182 9, function_name $L__info_string17, inlined_at 4 255 5
	mul.lo.s64 	%rd129, %rd10, 48;
	add.s64 	%rd130, %rd84, %rd129;
	add.s64 	%rd121, %rd130, 80;
	.loc	4 182 9, function_name $L__info_string17, inlined_at 4 255 5
	.loc	4 66 9, function_name $L__info_string18, inlined_at 4 182 9
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd120, %rd121;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r150,%r151,%r152,%r153}, [%rd120];
	// end inline asm
	.loc	4 183 9, function_name $L__info_string17, inlined_at 4 255 5
	add.s64 	%rd124, %rd130, 96;
	.loc	4 183 9, function_name $L__info_string17, inlined_at 4 255 5
	.loc	4 66 9, function_name $L__info_string18, inlined_at 4 183 9
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd123, %rd124;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r154,%r155,%r156,%r157}, [%rd123];
	// end inline asm
	.loc	4 184 9, function_name $L__info_string17, inlined_at 4 255 5
	add.s64 	%rd127, %rd130, 112;
	.loc	4 184 9, function_name $L__info_string17, inlined_at 4 255 5
	.loc	4 66 9, function_name $L__info_string18, inlined_at 4 184 9
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd126, %rd127;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r158,%r159,%r160,%r161}, [%rd126];
	// end inline asm
	bra.uni 	$L__BB0_14;

$L__BB0_4:
	.loc	4 310 10, function_name $L__info_string8, inlined_at 4 349 9
	setp.eq.s32 	%p4, %r14, 4;
	@%p4 bra 	$L__BB0_7;

	setp.ne.s32 	%p5, %r14, 1;
	@%p5 bra 	$L__BB0_14;

	.loc	4 321 13, function_name $L__info_string8, inlined_at 4 349 9
	.loc	3 1663 5, function_name $L__info_string19, inlined_at 4 321 13
	// begin inline asm
	call (%rd14), _optix_get_static_transform_from_handle, (%rd12);
	// end inline asm
	.loc	4 322 13, function_name $L__info_string8, inlined_at 4 349 9
	add.s64 	%rd131, %rd14, 64;
	bra.uni 	$L__BB0_8;

$L__BB0_10:
	.loc	4 303 13, function_name $L__info_string8, inlined_at 4 349 9
	.loc	3 1670 5, function_name $L__info_string20, inlined_at 4 303 13
	// begin inline asm
	call (%rd27), _optix_get_srt_motion_transform_from_handle, (%rd12);
	// end inline asm
	.loc	4 304 13, function_name $L__info_string8, inlined_at 4 349 9
	.loc	4 268 5, function_name $L__info_string21, inlined_at 4 304 13
	.loc	4 66 9, function_name $L__info_string22, inlined_at 4 268 5
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd29, %rd27;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r28,%r29,%r30,%r31}, [%rd29];
	// end inline asm
	.loc	4 66 9, function_name $L__info_string22, inlined_at 4 268 5
	add.s64 	%rd33, %rd27, 16;
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd32, %rd33;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r32,%r33,%r34,%r35}, [%rd32];
	// end inline asm
	.loc	4 66 9, function_name $L__info_string22, inlined_at 4 268 5
	add.s64 	%rd36, %rd27, 32;
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd35, %rd36;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r36,%r37,%r38,%r39}, [%rd35];
	// end inline asm
	.loc	4 66 9, function_name $L__info_string22, inlined_at 4 268 5
	add.s64 	%rd39, %rd27, 48;
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd38, %rd39;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r40,%r41,%r42,%r43}, [%rd38];
	// end inline asm
	.loc	4 66 9, function_name $L__info_string22, inlined_at 4 268 5
	add.s64 	%rd42, %rd27, 64;
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd41, %rd42;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r44,%r45,%r46,%r47}, [%rd41];
	// end inline asm
	.loc	4 66 9, function_name $L__info_string22, inlined_at 4 268 5
	add.s64 	%rd45, %rd27, 80;
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd44, %rd45;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r48,%r49,%r50,%r51}, [%rd44];
	// end inline asm
	.loc	4 66 9, function_name $L__info_string22, inlined_at 4 268 5
	add.s64 	%rd48, %rd27, 96;
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd47, %rd48;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r52,%r53,%r54,%r55}, [%rd47];
	// end inline asm
	.loc	4 66 9, function_name $L__info_string22, inlined_at 4 268 5
	add.s64 	%rd51, %rd27, 112;
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd50, %rd51;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r56,%r57,%r58,%r59}, [%rd50];
	// end inline asm
	.loc	4 66 9, function_name $L__info_string22, inlined_at 4 268 5
	add.s64 	%rd54, %rd27, 128;
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd53, %rd54;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r60,%r61,%r62,%r63}, [%rd53];
	// end inline asm
	.loc	4 66 9, function_name $L__info_string22, inlined_at 4 268 5
	add.s64 	%rd57, %rd27, 144;
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd56, %rd57;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r64,%r65,%r66,%r67}, [%rd56];
	// end inline asm
	.loc	4 67 5, function_name $L__info_string22, inlined_at 4 268 5
	mov.b32 	%f14, %r31;
	mov.b32 	%f15, %r32;
	.loc	4 268 5, function_name $L__info_string21, inlined_at 4 304 13
	.loc	4 221 5, function_name $L__info_string14, inlined_at 4 268 5
	and.b32  	%r84, %r30, 65535;
	add.s32 	%r85, %r84, -1;
	cvt.rn.f32.s32 	%f16, %r85;
	.loc	4 227 5, function_name $L__info_string14, inlined_at 4 268 5
	sub.ftz.f32 	%f17, %f13, %f14;
	sub.ftz.f32 	%f18, %f15, %f14;
	div.approx.ftz.f32 	%f19, %f17, %f18;
	mul.ftz.f32 	%f20, %f19, %f16;
	.loc	4 227 5, function_name $L__info_string14, inlined_at 4 268 5
	.loc	5 979 3, function_name $L__info_string15, inlined_at 4 227 5
	min.ftz.f32 	%f21, %f16, %f20;
	.loc	4 227 5, function_name $L__info_string14, inlined_at 4 268 5
	.loc	5 1107 3, function_name $L__info_string16, inlined_at 4 227 5
	mov.f32 	%f22, 0f00000000;
	max.ftz.f32 	%f23, %f22, %f21;
	.loc	4 230 5, function_name $L__info_string14, inlined_at 4 268 5
	setp.num.ftz.f32 	%p7, %f23, %f23;
	selp.f32 	%f24, %f23, 0f00000000, %p7;
	.loc	4 233 5, function_name $L__info_string14, inlined_at 4 268 5
	cvt.rmi.ftz.f32.f32 	%f25, %f24;
	add.ftz.f32 	%f26, %f16, 0fBF800000;
	min.ftz.f32 	%f27, %f25, %f26;
	.loc	4 235 5, function_name $L__info_string14, inlined_at 4 268 5
	sub.ftz.f32 	%f28, %f24, %f27;
	.loc	4 236 5, function_name $L__info_string14, inlined_at 4 268 5
	cvt.rzi.ftz.s32.f32 	%r86, %f27;
	.loc	4 271 5, function_name $L__info_string21, inlined_at 4 304 13
	mul.wide.s32 	%rd71, %r86, 64;
	add.s64 	%rd60, %rd36, %rd71;
	.loc	4 275 5, function_name $L__info_string21, inlined_at 4 304 13
	.loc	4 195 5, function_name $L__info_string23, inlined_at 4 275 5
	.loc	4 66 9, function_name $L__info_string18, inlined_at 4 195 5
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd59, %rd60;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r68,%r69,%r70,%r71}, [%rd59];
	// end inline asm
	.loc	4 196 5, function_name $L__info_string23, inlined_at 4 275 5
	add.s64 	%rd63, %rd60, 16;
	.loc	4 196 5, function_name $L__info_string23, inlined_at 4 275 5
	.loc	4 66 9, function_name $L__info_string18, inlined_at 4 196 5
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd62, %rd63;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r72,%r73,%r74,%r75}, [%rd62];
	// end inline asm
	.loc	4 197 5, function_name $L__info_string23, inlined_at 4 275 5
	add.s64 	%rd66, %rd60, 32;
	.loc	4 197 5, function_name $L__info_string23, inlined_at 4 275 5
	.loc	4 66 9, function_name $L__info_string18, inlined_at 4 197 5
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd65, %rd66;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r76,%r77,%r78,%r79}, [%rd65];
	// end inline asm
	.loc	4 198 5, function_name $L__info_string23, inlined_at 4 275 5
	add.s64 	%rd69, %rd60, 48;
	.loc	4 198 5, function_name $L__info_string23, inlined_at 4 275 5
	.loc	4 66 9, function_name $L__info_string18, inlined_at 4 198 5
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd68, %rd69;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r80,%r81,%r82,%r83}, [%rd68];
	// end inline asm
	.loc	4 201 5, function_name $L__info_string23, inlined_at 4 275 5
	setp.leu.ftz.f32 	%p8, %f28, 0f00000000;
	@%p8 bra 	$L__BB0_14;

	.loc	4 204 9, function_name $L__info_string23, inlined_at 4 275 5
	add.s64 	%rd73, %rd60, 64;
	.loc	4 204 9, function_name $L__info_string23, inlined_at 4 275 5
	.loc	4 66 9, function_name $L__info_string18, inlined_at 4 204 9
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd72, %rd73;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r87,%r88,%r89,%r90}, [%rd72];
	// end inline asm
	.loc	4 205 9, function_name $L__info_string23, inlined_at 4 275 5
	add.s64 	%rd76, %rd63, 64;
	.loc	4 205 9, function_name $L__info_string23, inlined_at 4 275 5
	.loc	4 66 9, function_name $L__info_string18, inlined_at 4 205 9
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd75, %rd76;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r91,%r92,%r93,%r94}, [%rd75];
	// end inline asm
	.loc	4 206 9, function_name $L__info_string23, inlined_at 4 275 5
	add.s64 	%rd79, %rd66, 64;
	.loc	4 206 9, function_name $L__info_string23, inlined_at 4 275 5
	.loc	4 66 9, function_name $L__info_string18, inlined_at 4 206 9
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd78, %rd79;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r95,%r96,%r97,%r98}, [%rd78];
	// end inline asm
	.loc	4 207 9, function_name $L__info_string23, inlined_at 4 275 5
	add.s64 	%rd82, %rd69, 64;
	.loc	4 207 9, function_name $L__info_string23, inlined_at 4 275 5
	.loc	4 66 9, function_name $L__info_string18, inlined_at 4 207 9
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd81, %rd82;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r99,%r100,%r101,%r102}, [%rd81];
	// end inline asm
	.loc	4 213 9, function_name $L__info_string23, inlined_at 4 275 5
	bra.uni 	$L__BB0_14;

$L__BB0_7:
	.loc	4 317 45, function_name $L__info_string8, inlined_at 4 349 9
	.loc	3 1705 5, function_name $L__info_string24, inlined_at 4 317 45
	// begin inline asm
	call (%rd131), _optix_get_instance_inverse_transform_from_handle, (%rd12);
	// end inline asm

$L__BB0_8:
	.loc	4 325 9, function_name $L__info_string8, inlined_at 4 349 9
	.loc	4 66 9, function_name $L__info_string18, inlined_at 4 325 9
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd18, %rd131;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r16,%r17,%r18,%r19}, [%rd18];
	// end inline asm
	.loc	4 326 9, function_name $L__info_string8, inlined_at 4 349 9
	add.s64 	%rd22, %rd131, 16;
	.loc	4 326 9, function_name $L__info_string8, inlined_at 4 349 9
	.loc	4 66 9, function_name $L__info_string18, inlined_at 4 326 9
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd21, %rd22;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r20,%r21,%r22,%r23}, [%rd21];
	// end inline asm
	.loc	4 327 9, function_name $L__info_string8, inlined_at 4 349 9
	add.s64 	%rd25, %rd131, 32;
	.loc	4 327 9, function_name $L__info_string8, inlined_at 4 349 9
	.loc	4 66 9, function_name $L__info_string18, inlined_at 4 327 9
	.loc	4 51 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	cvta.to.global.u64 %rd24, %rd25;
	// end inline asm
	.loc	4 53 5, function_name $L__info_string13, inlined_at 4 66 9
	// begin inline asm
	ld.global.v4.u32 {%r24,%r25,%r26,%r27}, [%rd24];
	// end inline asm

$L__BB0_14:
	.loc	4 344 40, function_name $L__info_string5, inlined_at 3 1635 5
	add.s32 	%r162, %r162, 1;
	.loc	4 344 5, function_name $L__info_string5, inlined_at 3 1635 5
	setp.lt.u32 	%p11, %r162, %r11;
	@%p11 bra 	$L__BB0_3;

$L__BB0_15:
	.loc	1 35 1
	ret;

}
	.file	1 "G:\\JadeSnakeSrc\\OptiXBaker\\optixBaker\\cuda\\test.cu"
	.file	2 "G:\\JadeSnakeSrc\\OptiXBaker\\optixBaker\\cuda\\Interaction.h"
	.file	3 "C:\\ProgramData\\NVIDIA Corporation\\OptiX SDK 8.0.0\\include\\internal/optix_device_impl.h"
	.file	4 "C:\\ProgramData\\NVIDIA Corporation\\OptiX SDK 8.0.0\\include\\internal/optix_device_impl_transformations.h"
	.file	5 "C:\\Program Files\\NVIDIA GPU Computing Toolkit\\CUDA\\v12.6\\include\\crt\\math_functions.hpp"
	.section	.debug_str
	{
$L__info_string0:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,54,103,101,116,80,82,68,73,49,57,80,101,114,73,110,116,101,114,115,101,99,116,105,111,110,68,97,116,97,69,69,80,84,95,118,0
$L__info_string1:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,49,55,111,112,116,105,120,71,101,116,80,97,121,108,111,97,100,95,48,69,118,0
$L__info_string2:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,49,55,111,112,116,105,120,71,101,116,80,97,121,108,111,97,100,95,49,69,118,0
$L__info_string3:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,52,50,111,112,116,105,120,84,114,97,110,115,102,111,114,109,78,111,114,109,97,108,70,114,111,109,79,98,106,101,99,116,84,111,87,111,114,108
.b8 100,83,112,97,99,101,69,54,102,108,111,97,116,51,0
$L__info_string4:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,50,53,111,112,116,105,120,71,101,116,84,114,97,110,115,102,111,114,109,76,105,115,116,83,105,122,101,69,118,0
$L__info_string5:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,49,48,111,112,116,105,120,95,105,109,112,108,51,54,111,112,116,105,120,71,101,116,87,111,114,108,100,84,111,79,98,106,101,99,116,84,114,97
.b8 110,115,102,111,114,109,77,97,116,114,105,120,69,82,54,102,108,111,97,116,52,83,50,95,83,50,95,0
$L__info_string6:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,49,53,111,112,116,105,120,71,101,116,82,97,121,84,105,109,101,69,118,0
$L__info_string7:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,50,55,111,112,116,105,120,71,101,116,84,114,97,110,115,102,111,114,109,76,105,115,116,72,97,110,100,108,101,69,106,0
$L__info_string8:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,49,48,111,112,116,105,120,95,105,109,112,108,52,52,111,112,116,105,120,71,101,116,73,110,116,101,114,112,111,108,97,116,101,100,84,114,97,110
.b8 115,102,111,114,109,97,116,105,111,110,70,114,111,109,72,97,110,100,108,101,69,82,54,102,108,111,97,116,52,83,50,95,83,50,95,121,102,98,0
$L__info_string9:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,51,49,111,112,116,105,120,71,101,116,84,114,97,110,115,102,111,114,109,84,121,112,101,70,114,111,109,72,97,110,100,108,101,69,121,0
$L__info_string10:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,51,57,111,112,116,105,120,71,101,116,77,97,116,114,105,120,77,111,116,105,111,110,84,114,97,110,115,102,111,114,109,70,114,111,109,72,97,110
.b8 100,108,101,69,121,0
$L__info_string11:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,49,48,111,112,116,105,120,95,105,109,112,108,51,52,111,112,116,105,120,71,101,116,73,110,116,101,114,112,111,108,97,116,101,100,84,114,97,110
.b8 115,102,111,114,109,97,116,105,111,110,69,82,54,102,108,111,97,116,52,83,50,95,83,50,95,80,75,50,54,79,112,116,105,120,77,97,116,114,105,120
.b8 77,111,116,105,111,110,84,114,97,110,115,102,111,114,109,102,0
$L__info_string12:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,49,48,111,112,116,105,120,95,105,109,112,108,50,52,111,112,116,105,120,76,111,97,100,82,101,97,100,79,110,108,121,65,108,105,103,110,49,54
.b8 73,50,54,79,112,116,105,120,77,97,116,114,105,120,77,111,116,105,111,110,84,114,97,110,115,102,111,114,109,69,69,84,95,80,75,83,51,95,0
$L__info_string13:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,49,48,111,112,116,105,120,95,105,109,112,108,56,111,112,116,105,120,76,100,103,69,121,0
$L__info_string14:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,49,48,111,112,116,105,120,95,105,109,112,108,50,49,111,112,116,105,120,82,101,115,111,108,118,101,77,111,116,105,111,110,75,101,121,69,82,102
.b8 82,105,82,75,49,56,79,112,116,105,120,77,111,116,105,111,110,79,112,116,105,111,110,115,102,0
$L__info_string15:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,51,109,105,110,69,102,102,0
$L__info_string16:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,51,109,97,120,69,102,102,0
$L__info_string17:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,49,48,111,112,116,105,120,95,105,109,112,108,51,48,111,112,116,105,120,76,111,97,100,73,110,116,101,114,112,111,108,97,116,101,100,77,97,116
.b8 114,105,120,75,101,121,69,82,54,102,108,111,97,116,52,83,50,95,83,50,95,80,75,83,49,95,102,0
$L__info_string18:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,49,48,111,112,116,105,120,95,105,109,112,108,50,52,111,112,116,105,120,76,111,97,100,82,101,97,100,79,110,108,121,65,108,105,103,110,49,54
.b8 73,54,102,108,111,97,116,52,69,69,84,95,80,75,83,51,95,0
$L__info_string19:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,51,51,111,112,116,105,120,71,101,116,83,116,97,116,105,99,84,114,97,110,115,102,111,114,109,70,114,111,109,72,97,110,100,108,101,69,121,0
$L__info_string20:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,51,54,111,112,116,105,120,71,101,116,83,82,84,77,111,116,105,111,110,84,114,97,110,115,102,111,114,109,70,114,111,109,72,97,110,100,108,101
.b8 69,121,0
$L__info_string21:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,49,48,111,112,116,105,120,95,105,109,112,108,51,52,111,112,116,105,120,71,101,116,73,110,116,101,114,112,111,108,97,116,101,100,84,114,97,110
.b8 115,102,111,114,109,97,116,105,111,110,69,82,54,102,108,111,97,116,52,83,50,95,83,50,95,80,75,50,51,79,112,116,105,120,83,82,84,77,111,116
.b8 105,111,110,84,114,97,110,115,102,111,114,109,102,0
$L__info_string22:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,49,48,111,112,116,105,120,95,105,109,112,108,50,52,111,112,116,105,120,76,111,97,100,82,101,97,100,79,110,108,121,65,108,105,103,110,49,54
.b8 73,50,51,79,112,116,105,120,83,82,84,77,111,116,105,111,110,84,114,97,110,115,102,111,114,109,69,69,84,95,80,75,83,51,95,0
$L__info_string23:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,49,48,111,112,116,105,120,95,105,109,112,108,50,55,111,112,116,105,120,76,111,97,100,73,110,116,101,114,112,111,108,97,116,101,100,83,114,116
.b8 75,101,121,69,82,54,102,108,111,97,116,52,83,50,95,83,50,95,83,50,95,80,75,83,49,95,102,0
$L__info_string24:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,52,50,111,112,116,105,120,71,101,116,73,110,115,116,97,110,99,101,73,110,118,101,114,115,101,84,114,97,110,115,102,111,114,109,70,114,111,109
.b8 72,97,110,100,108,101,69,121,0

	}

ERROR when optixModuleCreate:

[ 4][    COMPILER]: 
[ 2][    COMPILER]: COMPILE ERROR: No functions with semantic types found
Info: Module Statistics
	payload values        :          0
	attribute values      :          0
Info: Compiled Module Summary
	non-entry function(s):     1
	basic block(s)       :     1
	instruction(s)       :     2

What went wrong? I used the same CmakeLists.txt as optix_Sample and optix8.0 and cuda12.6，thanks！

[lspano]

Hi. It looks like no semantic function (e.g. closesthit, miss, …) is found. Are you sure the parameters to optixModuleCreate() are correct?

[dhart]

Hi @790989228,

It’s hard to say what’s wrong without a complete minimal reproducer. I see only a closesthit program in your source and PTX, but no raygen program. Do you have one? The raygen program is required.

–
David.

[790989228]

Now I’ve added a raygen program but still have a problem
“test.CU” content:

#include <crt/host_defines.h>
#include <optix_device.h>
#include "Interaction.h"
#include "LaunchParam.h"
extern "C" __constant__ OptixLaunchParams optixLaunchParams;

extern "C" __global__ void __closesthit__bsdf_radiance()
{
    const MeshSBTData &sbtData
      = *(const MeshSBTData*)optixGetSbtDataPointer();
    PerIntersectionData &prd = *getPRD<PerIntersectionData>();
    const int primID = optixGetPrimitiveIndex();
    int3 index  = sbtData.index_buffer[primID];
}

extern "C" __global__ void __raygen__rg()
{
    // Lookup our location within the launch grid
    const uint3 idx = optixGetLaunchIndex();
    const uint3 dim = optixGetLaunchDimensions();
}

createModule source:

optixContext = 0;
CUDA_CHECK( cudaFree( 0 ) );
CUcontext          cuCtx = 0;  // zero means take the current context
OPTIX_CHECK(optixInit());
OptixDeviceContextOptions options = {};
options.logCallbackFunction       = &context_log_cb;
options.logCallbackLevel          = 4;
OPTIX_CHECK( optixDeviceContextCreate( cuCtx, &options, &optixContext ) );
moduleCompileOptions = {};
moduleCompileOptions.maxRegisterCount = 50;
#if !defined(NDEBUG)
moduleCompileOptions.optLevel   = OPTIX_COMPILE_OPTIMIZATION_LEVEL_0;
moduleCompileOptions.debugLevel = OPTIX_COMPILE_DEBUG_LEVEL_FULL;
#endif
pipelineCompileOptions = {};
pipelineCompileOptions.traversableGraphFlags = OPTIX_TRAVERSABLE_GRAPH_FLAG_ALLOW_SINGLE_GAS;
pipelineCompileOptions.usesMotionBlur = false;
pipelineCompileOptions.numPayloadValues = 2;
pipelineCompileOptions.numAttributeValues = 2;
pipelineCompileOptions.exceptionFlags = OPTIX_EXCEPTION_FLAG_NONE;
pipelineCompileOptions.pipelineLaunchParamsVariableName = "optixLaunchParams";

pipelineLinkOptions.maxTraceDepth = 2;
size_t      inputSize = 0;
const char* input = sutil::getInputData( OPTIX_SAMPLE_NAME, OPTIX_SAMPLE_DIR, "test.cu", inputSize );
char log[2048];
size_t sizeof_log = sizeof(log);
OPTIX_CHECK_LOG(optixModuleCreate(optixContext,
    &moduleCompileOptions,
    &pipelineCompileOptions,
    input,
    inputSize,
    log,&sizeof_log,
    &module
));
if (sizeof_log > 1)
    PRINT(log);

sutil::getInputData got the ptx source:

//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-34431801
// Cuda compilation tools, release 12.6, V12.6.20
// Based on NVVM 7.0.1
//

.version 8.5
.target sm_50
.address_size 64

	// .globl	__closesthit__bsdf_radiance
.global .align 4 .b8 __nv_static_28__e1187107_7_test_cu_6c9dc3b4_identity[64] = {0, 0, 128, 63, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 128, 63, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 128, 63, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 128, 63};
.const .align 4 .f32 __nv_static_28__e1187107_7_test_cu_6c9dc3b4_Pi = 0f40490FDB;
.const .align 4 .f32 __nv_static_28__e1187107_7_test_cu_6c9dc3b4_InvPi = 0f3EA2F983;

.visible .entry __closesthit__bsdf_radiance()
{
	.reg .b32 	%r<6>;
	.reg .b64 	%rd<2>;
	.loc	1 7 0


	.loc	1 11 5
	.loc	2 61 2, function_name $L__info_string0, inlined_at 1 11 5
	.loc	3 1015 5, function_name $L__info_string1, inlined_at 2 61 2
	mov.u32 	%r2, 0;
	// begin inline asm
	call (%r1), _optix_get_payload, (%r2);
	// end inline asm
	.loc	2 62 2, function_name $L__info_string0, inlined_at 1 11 5
	.loc	3 1022 5, function_name $L__info_string2, inlined_at 2 62 2
	mov.u32 	%r4, 1;
	// begin inline asm
	call (%r3), _optix_get_payload, (%r4);
	// end inline asm
	.loc	1 14 1
	ret;

}
	// .globl	__raygen__rg
.visible .entry __raygen__rg()
{
	.reg .b32 	%r<7>;
	.loc	1 16 0


	.loc	1 21 1
	ret;

}
	.file	1 "G:\\JadeSnakeSrc\\OptiXBaker\\optixBaker\\cuda\\test.cu"
	.file	2 "G:\\JadeSnakeSrc\\OptiXBaker\\optixBaker\\cuda\\Interaction.h"
	.file	3 "C:\\ProgramData\\NVIDIA Corporation\\OptiX SDK 8.0.0\\include\\internal/optix_device_impl.h"
	.section	.debug_str
	{
$L__info_string0:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,54,103,101,116,80,82,68,73,49,57,80,101,114,73,110,116,101,114,115,101,99,116,105,111,110,68,97,116,97,69,69,80,84,95,118,0
$L__info_string1:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,49,55,111,112,116,105,120,71,101,116,80,97,121,108,111,97,100,95,48,69,118,0
$L__info_string2:
.b8 95,90,78,51,55,95,73,78,84,69,82,78,65,76,95,101,49,49,56,55,49,48,55,95,55,95,116,101,115,116,95,99,117,95,54,99,57,100,99,51
.b8 98,52,49,55,111,112,116,105,120,71,101,116,80,97,121,108,111,97,100,95,49,69,118,0

	}

thanks for help!

[790989228]

I solved the problem by using optixir instead of ptx