blob: 4a30aafaa611c3606c3e15190eb0e0eff1fc7a05 (
plain)
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
|
//TEST:SIMPLE(filecheck=CHECK_HLSL): -target hlsl -stage compute -entry computeMain
//TEST:SIMPLE(filecheck=CHECK_CUDA): -target cuda -stage compute -entry computeMain
//TEST_INPUT:ubuffer(data=[0 0 0 0 0 0 0 0 0 0 0 0 0 0 0 0], stride=4):out,name=outputBuffer
RWStructuredBuffer<int> outputBuffer;
// HLSL preserves WaveSize attribute in generated code
// CHECK_HLSL: [WaveSize(4)]
// CUDA doesn't emit WaveSize attribute but generates valid compute kernel
// CHECK_CUDA: RWStructuredBuffer<int> [[OUTPUTBUFFER:outputBuffer[_0-9]*]];
// CHECK_CUDA: extern "C" __global__ void computeMain()
// CHECK_CUDA: *(&([[GLOBALPARAMS:globalParams[_0-9]*]]->[[OUTPUTBUFFER]])[[[TID:tid[_0-9]*]]]) = [[TID]];
[WaveSize(4)]
[numthreads(4, 1, 1)]
void computeMain(int3 dispatchThreadID : SV_DispatchThreadID)
{
int tid = dispatchThreadID.x;
outputBuffer[tid] = tid;
}
|
