drbh HF Staff commited on
Commit
41a4f6d
·
verified ·
1 Parent(s): c148255

Upload folder using huggingface_hub

Browse files
Library/Logs/ScaleFT/sft/sft-2026-02-11T10_12_39-05_00.log DELETED
@@ -1,5 +0,0 @@
1
- 2026-02-11T10:12:39.276-0500 INFO sft command {"version": "1.98.1", "pid": 7488, "args": ["/usr/local/bin/sft", "resolve", "-q", "hf.co"], "log_directory": "Library/Logs/ScaleFT/sft"}
2
- 2026-02-11T10:12:39.282-0500 DEBUG macOS User Defaults check returned error {"error": "exit status 1"}
3
- 2026-02-11T10:12:39.283-0500 DEBUG Did not load a config from file, using default values {"path": "Library/Application Support/ScaleFT/sft.conf", "error": "stat Library/Application Support/ScaleFT/sft.conf: no such file or directory"}
4
- 2026-02-11T10:12:39.283-0500 ERROR CLI Action failed {"error": "", "action": "resolve"}
5
- 2026-02-11T10:12:39.283-0500 INFO RecordSpan {"t": "trace", "operation": "cli.resolve", "start": "2026-02-11T10:12:39.283-0500", "duration": "126.584µs", "traceID": 9173706215567653899, "spanID": 4785670708162327726, "tags": {"error":true}}
 
 
 
 
 
 
__KERNEL_NAME_NORMALIZED___metal/__KERNEL_NAME_NORMALIZED__.metal CHANGED
@@ -1,14 +1,8 @@
1
  #include <metal_stdlib>
2
  using namespace metal;
3
 
4
- kernel void __KERNEL_NAME_NORMALIZED___forward_kernel_float(device const float *input [[buffer(0)]],
5
- device float *output [[buffer(1)]],
6
- uint index [[thread_position_in_grid]]) {
7
  output[index] = input[index] + 1.0f;
8
  }
9
-
10
- kernel void __KERNEL_NAME_NORMALIZED___forward_kernel_half(device const half *input [[buffer(0)]],
11
- device half *output [[buffer(1)]],
12
- uint index [[thread_position_in_grid]]) {
13
- output[index] = input[index] + static_cast<half>(1.0);
14
- }
 
1
  #include <metal_stdlib>
2
  using namespace metal;
3
 
4
+ kernel void __KERNEL_NAME_NORMALIZED___kernel(device const float *input [[buffer(0)]],
5
+ device float *output [[buffer(1)]],
6
+ uint index [[thread_position_in_grid]]) {
7
  output[index] = input[index] + 1.0f;
8
  }
 
 
 
 
 
 
__KERNEL_NAME_NORMALIZED___metal/__KERNEL_NAME_NORMALIZED__.mm CHANGED
@@ -16,27 +16,27 @@ static inline id<MTLBuffer> getMTLBufferStorage(const torch::Tensor &tensor) {
16
  void __KERNEL_NAME_NORMALIZED__(torch::Tensor &out, torch::Tensor const &input) {
17
  TORCH_CHECK(input.device().is_mps(), "input must be a MPS tensor");
18
  TORCH_CHECK(input.is_contiguous(), "input must be contiguous");
19
- TORCH_CHECK(input.scalar_type() == torch::kFloat ||
20
- input.scalar_type() == torch::kHalf,
21
- "only float32 and float16 supported");
22
  TORCH_CHECK(input.sizes() == out.sizes(), "Tensors must have same shape");
23
- TORCH_CHECK(input.scalar_type() == out.scalar_type(), "Tensors must have same dtype");
24
- TORCH_CHECK(input.device() == out.device(), "Tensors must be on same device");
 
 
25
 
26
  @autoreleasepool {
27
  id<MTLDevice> device = MTLCreateSystemDefaultDevice();
28
  int numThreads = input.numel();
29
 
30
  NSError *error = nil;
31
- id<MTLLibrary> library = EMBEDDED_METALLIB_NAMESPACE::createLibrary(device, &error);
 
32
  TORCH_CHECK(library, "Failed to create Metal library: ",
33
  error.localizedDescription.UTF8String);
34
 
35
- std::string kernel_name = std::string("__KERNEL_NAME_NORMALIZED___forward_kernel_") +
36
- (input.scalar_type() == torch::kFloat ? "float" : "half");
37
- id<MTLFunction> func = [library newFunctionWithName:
38
- [NSString stringWithUTF8String:kernel_name.c_str()]];
39
- TORCH_CHECK(func, "Failed to create function: ", kernel_name.c_str());
40
 
41
  id<MTLComputePipelineState> pso =
42
  [device newComputePipelineStateWithFunction:func error:&error];
@@ -53,9 +53,10 @@ void __KERNEL_NAME_NORMALIZED__(torch::Tensor &out, torch::Tensor const &input)
53
  offset:out.storage_offset() * out.element_size()
54
  atIndex:1];
55
 
56
- NSUInteger tgSize = MIN(pso.maxTotalThreadsPerThreadgroup, (NSUInteger)numThreads);
 
57
  [encoder dispatchThreads:MTLSizeMake(numThreads, 1, 1)
58
- threadsPerThreadgroup:MTLSizeMake(tgSize, 1, 1)];
59
  [encoder endEncoding];
60
  torch::mps::commit();
61
  });
 
16
  void __KERNEL_NAME_NORMALIZED__(torch::Tensor &out, torch::Tensor const &input) {
17
  TORCH_CHECK(input.device().is_mps(), "input must be a MPS tensor");
18
  TORCH_CHECK(input.is_contiguous(), "input must be contiguous");
19
+ TORCH_CHECK(input.scalar_type() == at::ScalarType::Float,
20
+ "__KERNEL_NAME_NORMALIZED__ only supports float32");
 
21
  TORCH_CHECK(input.sizes() == out.sizes(), "Tensors must have same shape");
22
+ TORCH_CHECK(input.scalar_type() == out.scalar_type(),
23
+ "Tensors must have same dtype");
24
+ TORCH_CHECK(input.device() == out.device(),
25
+ "Tensors must be on same device");
26
 
27
  @autoreleasepool {
28
  id<MTLDevice> device = MTLCreateSystemDefaultDevice();
29
  int numThreads = input.numel();
30
 
31
  NSError *error = nil;
32
+ id<MTLLibrary> library =
33
+ EMBEDDED_METALLIB_NAMESPACE::createLibrary(device, &error);
34
  TORCH_CHECK(library, "Failed to create Metal library: ",
35
  error.localizedDescription.UTF8String);
36
 
37
+ id<MTLFunction> func =
38
+ [library newFunctionWithName:@"__KERNEL_NAME_NORMALIZED___kernel"];
39
+ TORCH_CHECK(func, "Failed to create function");
 
 
40
 
41
  id<MTLComputePipelineState> pso =
42
  [device newComputePipelineStateWithFunction:func error:&error];
 
53
  offset:out.storage_offset() * out.element_size()
54
  atIndex:1];
55
 
56
+ NSUInteger tgSize =
57
+ MIN(pso.maxTotalThreadsPerThreadgroup, (NSUInteger)numThreads);
58
  [encoder dispatchThreads:MTLSizeMake(numThreads, 1, 1)
59
+ threadsPerThreadgroup:MTLSizeMake(tgSize, 1, 1)];
60
  [encoder endEncoding];
61
  torch::mps::commit();
62
  });